7 ms·

What happens when you run a CUDA kernel?

294 points by mezark 3mo ago

fooblaster 3mo ago
The hardware has some open documentation. You don't actually need to read the kernel source to find some of the method documentation or qmd formats. See https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/compute/clcdc0qmd.h https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...
einpoklum 3mo ago
First - nice writeup which goes into a lot of nooks and crannies. That said, a lot of the user-space "voodoo" is gone if you don't go through CUDA's "runtime API". If you use the driver API, take your kernel source as a string and compile it with NVIDIA's run-time compiler, you'll have better visibility into a lot (not all) of what's going on. For the "raw" version of this, look at: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Introduction/vectorAdd_nvrtc https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int... but for a much more readable, and still fully transparent modern-C++ API version of the same, try this: https://github.com/eyalroz/cuda-api-wrappers/blob/master/examples/modified_cuda_samples/vectorAdd_nvrtc/vectorAdd_nvrtc.cpp https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa... that's a sample program for my CUDA API wrappers (header-only) library.
- mschuetz 3mo ago
  I like the driver API because it allows treating Cuda kernels like hot-reloadable shaders. It's fun to develop while being able to change the code at runtime.
  - einpoklum 3mo ago
    > I like the driver API because it allows treating Cuda kernels like hot-reloadable shaders. It is also much more friendly for library authors; and easier to wrap; and actually exposes a bunch of features the "runtime API" doesn't. The difficulty with it is that there just so many API calls; dozens of calls just for copying, for example. That was part of my motivation for writing my wrappers - making the supposedly "lower-level" API more accessible and intuitive than the supposedly "higher-level" API; and better integrated with the other libraries: NVTX, NVRTC, PTX compiler, fatbin library etc. > It's fun to develop while being able to change the code at runtime. It's also _the_ way to debug your kernels: If you don't load them dynamically, you have to recompile your application or kernel test harness every time you make a change to the kernel.
orliesaurus 3mo ago
There are companies whose whole job right now is to optimize kernels so that things run faster. I wonder if those companies are going to be dethroned by some sort of like open source library that can do that really well (I bet Nvidia could release it any day.).. or if they're going to thrive and be acquired by the big providers as a `moat` to speed up their infrerence.
- spmurrayzzz 3mo ago
  Near-term acquihires are certainly a likely bet I think. But given model progress on related benchmarks like kernelbench [1], I do think a set of more commoditized solutions is also inevitable. The caveat though is that each new gen of hardware often comes with brand new constraints/features that a given generation of models haven't seen before (e.g. tcgen05 in blackwell was OOD at one point). As the models start to generalize better, this might not be a showstopper, but still an issue at least currently. [1] https://kernelbench.com/ https://kernelbench.com/
  - xpct 3mo ago
    Thank you for sharing the link. It's fascinating that the models can only do 10-20% on the hard subset, and I wonder why that is so. The fact that they can only get 30-40% out of the fp8 GEMM seems unintuitive to me, I would've expected a convergence near ~80%.
    - spmurrayzzz 3mo ago
      I'm not entirely up to date with the latest batch, but I've reviewed some of the rollouts in the past and my sense is that the models are surprisingly good at getting correct custom kernels in the happy path, but still weak at sustained/shape-robust workloads. Having to deal with writing the full path from scratch compounded by weird memory layouts, odd sizes, routing, unpacking quantized weights, etc. is definitely challenging. Also, at least a portion of this you could argue is arbitrary and entirely scoped to the eval itself. The fp8 GEMM score could be low simply because one of the shapes is fairly skinny (i.e. not enough math work to keep the compute engine busy for a meaningful amount of time).
maxothex 3mo ago
[flagged]
mschuetz 3mo ago
That was an interesting read. Also enjoyed reading about the semaphores in the default stream. It's great that cuda implicitly handles syncing of commands for users and makes parallel commands optional and opt-in via streams, unlike Vulkan which completely unloads the full complexity of syncing to users right from the start.
kinow 3mo ago
I just finished a master's on HPC where I had to take some classes on CUDA, MPI+CUDA, OpenCL. Reading an article like this before the classes would have been a lot helpful! Especially the part just before and after "What does it mean for a warp to be eligible?".
- brcmthrowaway 3mo ago
  What university?
  - kinow 3mo ago
    It is a joint programme by Universities Santiago de Compostela and Coruña. I attended it from Barcelona, but they are in Galicia, Spain. https://www.usc.gal/en/studies/masters/engineering-and-architecture/master-high-performance-computing-online https://www.usc.gal/en/studies/masters/engineering-and-archi...
Jeeetendra 3mo ago
[flagged]
aberrahmane_b 3mo ago
It's very useful. The doorbell and QMD part were the most useful for me, because it connects the CUDA launch syntax to what actually gets submitted to the GPU. Most explanations stop around kernels, blocks and warps, but this made the CPU to driver to GPU path much easier to follow.
saagarjha 3mo ago
Control codes are a little more complicated than the post describes, they're really a table lookup rather than just bits in the control word.
effnorwood 3mo ago
on barra metal?
b112 3mo ago
Nice post. A note if the author is about, I had to use ublock origin to remove the header, as it hid text at the start of each page when printing. Firefox. (I prefer to read longer articles on my e-ink device via epub or PDF)
rampseeker 3mo ago
Wow, as soon as I opened it, there was even a detailed explanation of the vadd function I’d never seen before, but I couldn’t understand any of it.

294 points by mezark 3mo ago

The hardware has some open documentation. You don't actually need to read the kernel source to find some of the method documentation or qmd formats. See https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/compute/clcdc0qmd.h https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

einpoklum 3mo ago

First - nice writeup which goes into a lot of nooks and crannies. That said, a lot of the user-space "voodoo" is gone if you don't go through CUDA's "runtime API". If you use the driver API, take your kernel source as a string and compile it with NVIDIA's run-time compiler, you'll have better visibility into a lot (not all) of what's going on. For the "raw" version of this, look at: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Introduction/vectorAdd_nvrtc https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int... but for a much more readable, and still fully transparent modern-C++ API version of the same, try this: https://github.com/eyalroz/cuda-api-wrappers/blob/master/examples/modified_cuda_samples/vectorAdd_nvrtc/vectorAdd_nvrtc.cpp https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa... that's a sample program for my CUDA API wrappers (header-only) library.

mschuetz 3mo ago
I like the driver API because it allows treating Cuda kernels like hot-reloadable shaders. It's fun to develop while being able to change the code at runtime.
- einpoklum 3mo ago
  > I like the driver API because it allows treating Cuda kernels like hot-reloadable shaders. It is also much more friendly for library authors; and easier to wrap; and actually exposes a bunch of features the "runtime API" doesn't. The difficulty with it is that there just so many API calls; dozens of calls just for copying, for example. That was part of my motivation for writing my wrappers - making the supposedly "lower-level" API more accessible and intuitive than the supposedly "higher-level" API; and better integrated with the other libraries: NVTX, NVRTC, PTX compiler, fatbin library etc. > It's fun to develop while being able to change the code at runtime. It's also _the_ way to debug your kernels: If you don't load them dynamically, you have to recompile your application or kernel test harness every time you make a change to the kernel.

orliesaurus 3mo ago

There are companies whose whole job right now is to optimize kernels so that things run faster. I wonder if those companies are going to be dethroned by some sort of like open source library that can do that really well (I bet Nvidia could release it any day.).. or if they're going to thrive and be acquired by the big providers as a `moat` to speed up their infrerence.

spmurrayzzz 3mo ago
Near-term acquihires are certainly a likely bet I think. But given model progress on related benchmarks like kernelbench [1], I do think a set of more commoditized solutions is also inevitable. The caveat though is that each new gen of hardware often comes with brand new constraints/features that a given generation of models haven't seen before (e.g. tcgen05 in blackwell was OOD at one point). As the models start to generalize better, this might not be a showstopper, but still an issue at least currently. [1] https://kernelbench.com/ https://kernelbench.com/
- xpct 3mo ago
  Thank you for sharing the link. It's fascinating that the models can only do 10-20% on the hard subset, and I wonder why that is so. The fact that they can only get 30-40% out of the fp8 GEMM seems unintuitive to me, I would've expected a convergence near ~80%.
  - spmurrayzzz 3mo ago
    I'm not entirely up to date with the latest batch, but I've reviewed some of the rollouts in the past and my sense is that the models are surprisingly good at getting correct custom kernels in the happy path, but still weak at sustained/shape-robust workloads. Having to deal with writing the full path from scratch compounded by weird memory layouts, odd sizes, routing, unpacking quantized weights, etc. is definitely challenging. Also, at least a portion of this you could argue is arbitrary and entirely scoped to the eval itself. The fp8 GEMM score could be low simply because one of the shapes is fairly skinny (i.e. not enough math work to keep the compute engine busy for a meaningful amount of time).

maxothex 3mo ago

[flagged]

mschuetz 3mo ago

That was an interesting read. Also enjoyed reading about the semaphores in the default stream. It's great that cuda implicitly handles syncing of commands for users and makes parallel commands optional and opt-in via streams, unlike Vulkan which completely unloads the full complexity of syncing to users right from the start.

kinow 3mo ago

I just finished a master's on HPC where I had to take some classes on CUDA, MPI+CUDA, OpenCL. Reading an article like this before the classes would have been a lot helpful! Especially the part just before and after "What does it mean for a warp to be eligible?".

brcmthrowaway 3mo ago
What university?
- kinow 3mo ago
  It is a joint programme by Universities Santiago de Compostela and Coruña. I attended it from Barcelona, but they are in Galicia, Spain. https://www.usc.gal/en/studies/masters/engineering-and-architecture/master-high-performance-computing-online https://www.usc.gal/en/studies/masters/engineering-and-archi...

Jeeetendra 3mo ago

[flagged]

aberrahmane_b 3mo ago

It's very useful. The doorbell and QMD part were the most useful for me, because it connects the CUDA launch syntax to what actually gets submitted to the GPU. Most explanations stop around kernels, blocks and warps, but this made the CPU to driver to GPU path much easier to follow.

saagarjha 3mo ago

Control codes are a little more complicated than the post describes, they're really a table lookup rather than just bits in the control word.

effnorwood 3mo ago

on barra metal?

b112 3mo ago

Nice post. A note if the author is about, I had to use ublock origin to remove the header, as it hid text at the start of each page when printing. Firefox. (I prefer to read longer articles on my e-ink device via epub or PDF)

rampseeker 3mo ago

Wow, as soon as I opened it, there was even a detailed explanation of the vadd function I’d never seen before, but I couldn’t understand any of it.