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

CUDA Books

github.com

Post reply on HN
Print view View original

51–60 of 63 posts

Re: CUDA Books

#51

Post by namibj »

fransje26 wrote: post #49
jpgvm wrote: post #45

Earlier quoted context omitted.

Yeah pretty much this. I would separate the knowledge into maybe 3 distinct buckets. The baseline: device/host boundary, SIMT programming etc. The intermediate: kernel architecture, CUDA graph vs persistent kernels, warp specialisation/divergence avoidance techniques etc. The advanced: architecture specifics so tcgen05, TMA, SMEM/HBM, memory throughput vs compute biases in various arch impls., GEMM, FHMA, all the tri…

> [..] all the tricks that make modern fused kernels very fast This would require very different (re-written?) kernels than a few years back, wouldn't it? Would you have any good resources on the topic?

There's actually little that changed in a way too fundamentally to matter other than _perhaps_ getting the async load-from-global-to-shared-memory DMA memcpy that avoided blocking register file space as target buffers for in-flight read-from-global operations. Shared after all is just a partition of L1d$ since iirc Volta (since they offered non-fixed/at-launch-requested expanded shared capacity support), so it made sense to provide this not-just-a-hint "prefetch into this user-managed slice of what is otherwise L1d$": it's AFAIK basically just some special load-like units that ask special L1d$-miss-fill units to deliver to a now-explicitly-specified target location in the non-automatic-cache partition of the local SRAM and signal completion in otherwise fairly normal local semaphore/barrier fashion.

The major difference is that this doesn't have a natural moment to transform/touch the values after read from global and before storage to shared.

Otherwise, tiled MMA (gemm) kernels where normal even in Maxwell days (after the classic K80, before the P100; Maxwell is when H.265 support landed).

Re: CUDA Books

#52

Post by KeplerBoy »

boomzilla wrote: post #36
somethingsome wrote: post #14

Having read or at least skimmed most of those books, I think the best intro is 'CUDA Programming: A Developer's Guide to Parallel Computing with GPUs' Massively Parallel Processors: A Hands-on Approach is not really good in my opinion, many small mistakes and confusing sentences (even when you know cuda). CUDA by Example: An Introduction to General-Purpose GPU Programming is too simple and abstract too much the archi…

How about this guide: https://docs.nvidia.com/cuda/cuda-programming-guide/pdf/cuda...

That's hardly a guide. It's the defacto documentation, you have to read this either way.

Re: CUDA Books

#53

Post by KeplerBoy »

jpgvm wrote: post #45
somethingsome wrote: post #21

Earlier quoted context omitted.

Not really, Hardware didn't really change that much, of course you'll not find Tensor or raytracing cores, but you will have a very solid grasp of gpu programming and the cuda language (that didn't change that much either), and then you can easily learn those more modern things with blog posts or even, at worst, chatgpt.

Yeah pretty much this. I would separate the knowledge into maybe 3 distinct buckets. The baseline: device/host boundary, SIMT programming etc. The intermediate: kernel architecture, CUDA graph vs persistent kernels, warp specialisation/divergence avoidance techniques etc. The advanced: architecture specifics so tcgen05, TMA, SMEM/HBM, memory throughput vs compute biases in various arch impls., GEMM, FHMA, all the tri…

I wish there were any good literature on GPU Direct RDMA and GPU NetIO. Got any tips?

Re: CUDA Books

#54

Post by cold_harbor »

for LLM work, reading the Flash Attention and vLLM kernel source taught me more than any book. real code makes memory hierarchy concrete — books stay too abstract.

Re: CUDA Books

#55

Post by jpgvm »

KeplerBoy wrote: post #53
jpgvm wrote: post #45

Earlier quoted context omitted.

Yeah pretty much this. I would separate the knowledge into maybe 3 distinct buckets. The baseline: device/host boundary, SIMT programming etc. The intermediate: kernel architecture, CUDA graph vs persistent kernels, warp specialisation/divergence avoidance techniques etc. The advanced: architecture specifics so tcgen05, TMA, SMEM/HBM, memory throughput vs compute biases in various arch impls., GEMM, FHMA, all the tri…

I wish there were any good literature on GPU Direct RDMA and GPU NetIO. Got any tips?

So I would say the most important thing is that the APIs these are using as in mlx5 DevX (essentially direct fw access) or ibverbs are exactly the same regardless if it's CPU or GPU talking to it. So with that in mind the source of rdma-core, DPDK, ucx etc may be the most elucidating when it comes to low level details.

For higher level patterns again the APIs are the same so anything building on libibverbs or aforementioned ucx etc are pretty compatible from a high level ideas perspective. If you are new to RDMA in general definitely start with raw verbs instead of using abstractions like MPI if you really want to build a good intuition and then move to MPI once you understand what it is doing for you.

Re: CUDA Books

#57

Post by adrian_b »

aaqaishtyaq wrote: post #56

I really need to buy nvidia GPUs to be able to learn CUDA.

Fortunately, unlike with the AMD GPUs, using one of the cheaper NVIDIA GPUs is sufficient for learning CUDA, because CUDA works similarly on all models.

An expensive NVIDIA GPU is required only if your purpose is not just to learn, but to actually do useful graphics or ML/AI work.

Re: CUDA Books

#58

Post by dahart »

tirutiru wrote: post #38
dahart wrote: post #13

Regarding the section on Python and high-level CUDA, anyone interested should maybe first take a peek at Warp, which I’m guessing is too new to have a book yet. Warp lets you write CUDA kernels directly in Python, and it’s a breeze to get started. https://github.com/nvidia/warp

It's a bit confusing now with Numba Cuda also being officially maintained by Nvidia. Also Cuda Python, which looks older. Which of these - warp, numba, cp, is the best bet for a beginner? https://nvidia.github.io/numba-cuda/ https://developer.nvidia.com/cuda/python

I haven’t tried them all, but I suspect Warp is the easiest; it’s ridiculously easy. I’m sure there are some tradeoffs, so once you learn a little CUDA in Python it might make sense to switch from Warp to Numba or CP depending on what you’re doing.

Re: CUDA Books

#59

Post by dahart »

saagarjha wrote: post #47
nnevatie wrote: post #42

Earlier quoted context omitted.

To do what? If you need the highest performance GPU kernel performance on NVidia HW, using CUDA is the way to go.

Writing efficient CUDA code is very, very difficult; most CUDA code is not actually good at utilizing the hardware. It is much easier to write performant code in higher level languages (and most people are doing exactly this).

That all depends on what you’re doing. Like I said, if a high level lang or lib supports and fits your goal well, then yes you should use it. I don’t know what most people are doing, but it’s fair to say that a lot of people can use a higher level language.

If you’re trying to learn CUDA, then using a higher level language is not the best approach. If you already used a high level language and found that your performance is lacking and could be better if you could fuse some of your kernels, and avoid some of the memory round-trips, then moving to something lower level is called for.

I’m suggesting it’s better to think about your goals for one minute and understand the basic choices than it is to assume there’s something that works for everyone’s goals, and higher level languages don’t meet everyone’s goals.

Re: CUDA Books

#60

Post by dandanua »

dahart wrote: post #13

Regarding the section on Python and high-level CUDA, anyone interested should maybe first take a peek at Warp, which I’m guessing is too new to have a book yet. Warp lets you write CUDA kernels directly in Python, and it’s a breeze to get started. https://github.com/nvidia/warp

You can also write CUDA kernels directly in Julia using CUDA.jl. I basically learned CUDA programming by experimenting in Julia with the help of LLMs.
Post reply on HN