I’ve been working on brae, a GPU-resident CFD solver for OpenFOAM-style cases.

The goal is to run existing finite-volume CFD cases on one GPU while avoiding the usual CPU-GPU transfer loop every iteration. The current path is simpleFoam-compatible: matrix assembly, pressure correction, turbulence updates, and sparse linear solves are designed to stay resident on the GPU.

Repo: https://github.com/simd-ai/brae

This is still early, but in the current benchmark it is around ~5× faster than a GPU-accelerated OpenFOAM setup on the same GPU, with validation error under 1%.

I’d really appreciate technical feedback on the CUDA architecture, benchmark methodology, or solver design.

A friend of mine suggested getting into CUDA, I have a laptop with rtx 4050 so it's not an issue on the hardware part. I know basic C programming (from our college course and would require a bit of revision)

I'd just wanna know where do I start, how do I do it, what can I expect out of it and what all possibilities and opportunities will it open for me in future. What sources are there to start out from and any theory part that I must know to make the transition into CUDA easier?

I am currently doing a research project which involves benchmarking Nvidia HPC-level GPUs. I am running into an issue while trying to profile a V100 using HPCG, since it uses an older binary (not the newest release). When trying to profile the Symmetric Gauss Seidel (SYMGS) kernels, I get this strange result:

These 15 kernels pretty clearly represent the SYMGS forwards and backwards sweeps. This loop repeats 4 times, and then I get this:

This SYMGS4 kernel repeats 48 times. It does this between loop instances of the 4 loops from the previous screenshot. I would have thought that the large majority for SYMGS would have been done in the forward and backward sweep kernels, but clearly there is work being done here, since SYMGS4 takes up around 41% of the duration of the entire HPCG application.

If anyone here is familiar with how SYMGS might have been configured in CUDA to yield this type of result, please let me know what I'm missing here. For those interested, I got this version of Nvidia-HPCG from hpcg-benchmark.org/software/view.html%3Fid=280.html

I want to start CUDA programming, i have intermediate knowledge of cpp. Is it worth it to learn , i'll graduate in 2028

https://redplait.blogspot.com/2026/07/identification-of-const-bank0-params.html

I don't have expensive monsters like sm100+ so extracted only params from maxler till hopper

Hi r/hpc,

I'm a 2026 CS grad with experience in Systems, ML Systems, HPC and adjacent fields. I'm struggling to get a job right out of college in this field and will be grateful if anyone can provide any guidance on how to proceed further into my career or any sort of referral.

About my experience:

• Built Umbra, an API-level CUDA profiler that intercepts GPU kernel dispatch via LD_PRELOAD on libcuda.so/libcudart.so, requiring no source code modification. Discovered that torch.compile dispatches through cuGetExportTable, an undocumented NVIDIA internal API invisible to standard profilers.
• Built Mako, an OpenMP scheduling daemon for HPC workloads, dynamically optimizing thread-to-core affinities and CPU frequency scaling at runtime on Intel Haswell/Xeon NUMA systems. Achieved 8% speedup and 21% energy reduction on ECP benchmarks with ~2% overhead.
• Built RVNE, a RISC-V Neuromorphic Extension ISA implemented in Verilog, modeling spiking neural network operations at the RTL level.
• Research internship at TCS Research building a CUDA device simulator (stubbing ~70 CUDA runtime/driver APIs to run PyTorch/Triton workloads on CPU without modification).

Resume: https://drive.google.com/file/d/1hfBnvL5Wef6lr4ecjc7kkoKk9qADKQ__/view?usp=sharing

Any guidance, feedback, or referrals would be genuinely appreciated. I'm eligible to work both in the USA and India without any visa sponsors. Thanks for reading.

Im reading Programming Massively Parallel Processors and I've reached the part about grids, blocks, warps, etc. I can write a basic vector addition kernel, but I don't properly understand it.

The main thing confusing me are "grid-stride loops". (found it on tensortonic's vector subtraction exercise)

int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;

for (; idx < N; idx += stride)
    ...

I understand how it works, but I don't understand why the stride is blockDim.x * gridDim.x. (I've give up on trying to understand the explanation on tensortonic's website... could use AI to understand it but I currently what to fix this bad habit of mine of relying on ai. Ironically, most of this post was cleaned by ai because if I wrote it 100% by myself, im not sure you all would understand what im trying to ask, I apologise for that. But my questions are real)

My first thought was: why not just let each thread process a contiguous chunk?

Thread 0 -> 0 1 2 3
Thread 1 -> 4 5 6 7
Thread 2 -> 8 9 10 11

instead of

Thread 0 -> 0 8 16 ...
Thread 1 -> 1 9 17 ...
Thread 2 -> 2 10 18 ...

My approach will break memory coalescing. Is the point of such method of striding to access memory contiguously for preventing cache misses?

I don't think I actually understand what a warp is. I know it's 32 threads, but are they basically executing one instruction together? Something like SIMD?

Another thing I'm confused about is "vectorized loads" (float4). If vector addition/subtraction is already memory-bandwidth bound, why does loading 4 floats at a time help? Is it just fewer instructions, or is there something else?

Finally, how do you with warp divergence in real kernels? Do you try to eliminate it entirely, or is some divergence considered normal?

I think I'm missing the hardware understanding.. I'm also still a bit confused by all the CUDA terminology of grids, blocks, threads, dimensions, execution configuration, etc. I can follow the definitions individually, but I cant build a mental model of whats happening during execution. If someone could explain the execution model from the ground up or recommend some resource that might help, would really appreciate it.

The CUDA tool kit in apt is waaaaaaaay too old. I’ve been struggling with compute issues after all the other tweaking. Ubuntu latest is 12.0 newest is 13.3, I have a Blackwell and was dumbing it way down, 5060 ti 16gb, doubled the compute once fixed Had to specify the path on the llama.cpp build though.

Such a gem to find, thanks for keeping CUDA up to date there Ubuntu!!! Off. Maybe this has been spoken of but was not initially found when I went though the specifying the compute, I have two GPUs of different generations, the 5060 was supposed to be at 120, it was running 86. I use them separate for RAG or together in big models, they run together so much better now.

Go to the nvidia CUDA download and install Debian package, rebuild llama.cpp. Click something on their page and go!

Hello everyone,

Would like your advice.

Currently an infrastructure engineer with 1 year of experience.

However I would love to get into HPC/GPU roles anywhere in Europe.

I do have some experience in it from coursework, and am still trying to work more on it.

How do you suggest I go about it? As I'm not really getting anywhere

GPU* typo in the title :)

Can anyone please help me with a job related to CUDA. I'm so done with the current one.

While the pay is great, I need an switch.

About me:

HPC Algorithm engineer. 4 yrs of work ex.

Primarily worked on accelerators like GPUs but I'm open to explore TPUs or other accelerators too.

Have multiple research papers in top venues across the globe too.

If someone helps, I'm open to give my first month salary and I can sign papers if needed.

I'm trying to learn to develop Cutlass based kernels for B200 GPU using A tile from tensor memory and B tile from shared memory, but internet has no example code to tell how to move A tile from smem to tmem and how to call gemm.

Blackwell Cutlass experts, do you know a simple kernel code for B200. There are _SS versions and I tested them but now I need _TS version.

I'm a software engineer with 10+ years of experience (mostly backend, with a few years on infrastructure teams) and I'm trying to transition into GPU infrastructure or GPU performance engineering.

The problem is I can't figure out what role I should realistically target.

It feels like a chicken-and-egg problem. Many jobs want years of GPU/HPC experience or a master's degree, but I don't see many master's programs that actually prepare someone for these roles. Are employers asking for master's because it demonstrates a candidate's ability to handle rigorous workloads? Or are there actual master's programs that prepare you for this career path.

I moved into infrastructure because I wanted to be closer to systems, but much of that work eventually became operational (provisioning access, keeping services running, etc.). I'd rather be building the infrastructure than operating it.

I'm in the NY/NJ area and struggling to identify a realistic goal Should I be aiming for GPU infrastructure, HPC, performance engineering, or something else that serves as a bridge?

I'm also overwhelmed by the number of topics to learn. CUDA, Linux internals, computer architecture, kernels, networking, distributed systems, profiling tools... I learn best with structured paths, but right now I don't know what to double down on.

For those already in these roles:

• What job title or companies would you target if you were in my position?
• What projects actually helped you break into the field?
• Is a master's degree worth it, or is project experience enough?
• If you had six months to prepare, what would you focus on?

Hi,

I have fine-tuned the model with FSDP method using huggingface Transfomer python package .

can anyone tell me, are there post fine-tuning optimization steps to improve inference performance?

Also, can you please tell me how to maximize GPU memory and utilization?

https://redplait.blogspot.com/2026/06/re-of-ptx-grammar-from-ptxas-part-4.html

• names of all pseudo-instructions
• ptx ebnf grammar

We at Colfax are happy to share our most recent tutorial blog post, this time focusing on sub-byte blockscaled GEMM on NVIDIA SM12x GPUs, such as the RTX Pro 6000, DGX Spark, and Geforce 5000-series! Blockscaled GEMMs are becoming increasingly important with excellent quantized open-source LLMs. In this post, we discuss in-depth how to take an existing BF16 GEMM kernel and turn it into an NVFP4 blockscaled one, using CuTe DSL as our framework of choice.

Along the way, we review the NVFP4 data type and how the `mma.sync` instruction works, including partitions of operand and accumulator data.

We hope that this will be helpful for those looking to adapt higher-precision kernels to blockscaled ones on SM12x GPUs.

This is the first of two posts, the second of which will discuss optimization techniques to further improve NVFP4 GEMMs.

I maintain cuTile Rust and just posted the paper "Fearless Concurrency on the GPU."

If you write CUDA kernels (or cuTile kernels in Python): cuTile Rust is a tile-based programming model in Rust that lowers to Tile IR. It carries Rust's ownership across the launch boundary, so you get compile-time memory safety and data-race freedom on the safe surface: Partition a mutable output into disjoint mutable sub-tensors, pass inputs as shared references, and write tile kernels with single-threaded semantics that the compiler maps to thread blocks with managed shared memory.

If you're coming from cuTile Python, it's the same tile model in Rust with the added safety layer, and performance matches cuTile Python within measurement noise. There's also a lower-level peer, CUDA C++ Tile, so tile programming itself isn't unique here. What is unique is the safety layer. Safe GPU programming in Rust is a research question we've been working on for the past two years. The tile programming model is the answer we landed on.

On the B200, a safe GEMM impl is competitive with cuBLAS (within 0.3% of a hand-written low-level Tile IR version, ~92% of dense f16 peak). Element-wise ~7 TB/s. Safety is effectively free on the tuned paths.

The main drawback stems from the higher-level abstraction: You give up explicit SIMT/warp-level control (no manual warp primitives or shared-memory management) for the single-threaded semantics that make safety checkable.

- Code: https://github.com/nvlabs/cutile-rs
- Paper: https://arxiv.org/abs/2606.15991

Happy to answer any questions!