Hacker News Board Hacker News, browsed like a classic message board
Live data from Hacker News

What happens when you run a CUDA kernel?

fergusfinn.com

Post reply on HN
Print view View original

11–20 of 36 posts

Re: What happens when you run a CUDA kernel?

#11

Post by connicpu »

orliesaurus wrote: post #5

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.

When you run CUDA at scale dealing with nvidia driver and library bugs takes up a disgustingly large percentage of engineer time, I don't know a lot of people who would be looking forward to rely on more nvidia libraries.

Re: What happens when you run a CUDA kernel?

#12

Post by einpoklum »

mschuetz wrote: post #4
einpoklum wrote: post #3

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-sam…

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.

> 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.

Re: What happens when you run a CUDA kernel?

#13

Post by einpoklum »

orliesaurus wrote: post #5

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.

Probably not, because the specifics of the workload - exact parameters, representation of data in memory, value ranges etc - lead you to highly divergent optimization strategies.

Re: What happens when you run a CUDA kernel?

#14

Post by orliesaurus »

connicpu wrote: post #11
orliesaurus wrote: post #5

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.

When you run CUDA at scale dealing with nvidia driver and library bugs takes up a disgustingly large percentage of engineer time, I don't know a lot of people who would be looking forward to rely on more nvidia libraries.

fair point, but are there alternatives that aren't CUDA locked?

Re: What happens when you run a CUDA kernel?

#15

Post by orliesaurus »

einpoklum wrote: post #13
orliesaurus wrote: post #5

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.

Probably not, because the specifics of the workload - exact parameters, representation of data in memory, value ranges etc - lead you to highly divergent optimization strategies.

shouldn't it be possible to be run as a mlautoresearch project? i.e. orchestrate 10 strategies to speed it up, run in paralellel, pick the winning and go from there?

Re: What happens when you run a CUDA kernel?

#16

Post by whattheheckheck »

connicpu wrote: post #11
orliesaurus wrote: post #5

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.

When you run CUDA at scale dealing with nvidia driver and library bugs takes up a disgustingly large percentage of engineer time, I don't know a lot of people who would be looking forward to rely on more nvidia libraries.

Is there an issue board for these bugs? I want to see what is a disgustingly large percent. 50%?

Re: What happens when you run a CUDA kernel?

#17

Post by einpoklum »

orliesaurus wrote: post #15
einpoklum wrote: post #13

Earlier quoted context omitted.

Probably not, because the specifics of the workload - exact parameters, representation of data in memory, value ranges etc - lead you to highly divergent optimization strategies.

shouldn't it be possible to be run as a mlautoresearch project? i.e. orchestrate 10 strategies to speed it up, run in paralellel, pick the winning and go from there?

You are assuming all problems in the world are solvable by one of "10 strategies".

Re: What happens when you run a CUDA kernel?

#18

Post by xpct »

spmurrayzzz wrote: post #7
orliesaurus wrote: post #5

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.

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…

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%.

Re: What happens when you run a CUDA kernel?

#19

Post by david-gpu »

connicpu wrote: post #11
orliesaurus wrote: post #5

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.

When you run CUDA at scale dealing with nvidia driver and library bugs takes up a disgustingly large percentage of engineer time, I don't know a lot of people who would be looking forward to rely on more nvidia libraries.

How do you determine that the bugs you run into are located in the Nvidia drivers and libraries?

Way back when I wrote the OpenCL driver at Qualcomm, we would frequently get bug reports from customers complaining about our code. During my tenure, every single one of them was root-caused as an application bug. Unsurprisingly, considering that our code was backed by an extensive test suite and their code wasn't.

Not to say that our code was perfect, of course. But people have a tendency to blame GPU drivers when the problem often lies elsewhere.

Re: What happens when you run a CUDA kernel?

#20

Post by connicpu »

david-gpu wrote: post #19
connicpu wrote: post #11

Earlier quoted context omitted.

When you run CUDA at scale dealing with nvidia driver and library bugs takes up a disgustingly large percentage of engineer time, I don't know a lot of people who would be looking forward to rely on more nvidia libraries.

How do you determine that the bugs you run into are located in the Nvidia drivers and libraries? Way back when I wrote the OpenCL driver at Qualcomm, we would frequently get bug reports from customers complaining about our code. During my tenure, every single one of them was root-caused as an application bug. Unsurprisingly, considering that our code was backed by an extensive test suite and their code wasn't. Not to…

If you're big enough you can get direct access to Nvidia engineers, and they are usually transparent when they find out the bug was in their software and send you a patched version to try to resolve the issue
Post reply on HN