CUDA Books | Hacker News Reader

CUDA Books

github.com

243 pointsdariubs60 comments

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

Next year I'm planning to start writing a cuda book that starts by engineering the hardware, and goes up to the optimization part on that harware (which is basically a nvidia card) including all the main algorithms (except for graphs).

I'm already teaching the course in this way at uni, and it is quite successful among students.

the first book was published in 2012,is it too outdated?

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 tricks that make modern fused kernels very fast. Also would bucket most GPU Direct RDMA/GPU NetIO/friends here too.

The baseline hasn't changed much and probably won't, the intermediate knowledge has also remained pretty reliably stable for ~10 years with only things like graphs changing stuff. Tile might become more relevant than it is today but for now CUDA, cuBLAS, friends are where it's worth investing knowledge.

> [..] 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).

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.

Thanks I appreciate the answer!

I really wish there were better options to PMPP... It's by far the most up-to-date book, but I totally agree the writing is sort of bad and some of the code examples are straight up incorrect.

So tl;dr, you have at least one person who would pay for a better book :-)

Very valuable comment. Thank you.

I always appreciate book lists like this one, but having a small targeted list is more practical for those of us with limited reading time.

Interesting, thanks for sharing.

What makes CUDA Programming: A Developer's Guide to Parallel Computing with GPUs better among its peers?

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.

Thank you, that is very useful advice !

This is highly condensed video of all important concepts in CUDA from Stephen Jones, one of the CUDA architects: https://www.youtube.com/watch?v=QQceTDjA4f4

Understand everything he talks about and you understand CUDA.

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.

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.

Increasingly (for instance ADSP podcast [1]) those in nvidia's inner circle are advocating against writing your own CUDA kernels. (Unless that's your full time job at nvidia, that is).

[1] https://adspthepodcast.com/2024/08/30/Episode-197.html

It’s not about whether you work at Nvidia. Avoid writing CUDA kernels if there are higher level libraries that do what you need. Do write CUDA kernels if you want to learn how, or if you need the low level control, or to micro-optimize. Being able to fuse kernels to avoid memory traffic or get better specialization is also a reason to reach for raw CUDA. Just consider what’s the right tool for the job…

I don't think writing CUDA is a good way to do this tbh

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.

I think there are very few things that should be written in CUDA and many of them are just people who like to write CUDA for the fun of it

That would be cool but nvidia released blackwell and still have not released unbroken kernels for sm120. Sm120 is not the data center gpu, so it doesn't get its love. So we can't depend on nvidia to do the right thing is my point unfortunately

That advice seems like nonsense. It's like saying avoid C because you can use Python, or avoid writing a graphics engine because you can license Unreal.

Not at all, the advice is like use SDL or Raylib instead of writing your framebuffer blitter in inline Assembly to call from C.

I bet you will learn alot doing that though

Depends if the purpose is learning or actually delivery something on the same amount of time.

Each one has their place.

can very much agree about not writing stuff like reductions yourself, unless you have good reason to. but this sort of feels like another "implement everything with <nvidia stuff> and you'll have a great time!! (but also coincidentally get locked in even more to Nvidia hardware)"

"AI Systems Performance Engineering" might deserve a mention, even though it's not strictly CUDA.

I liked going through https://www.olcf.ornl.gov/cuda-training-series/ for an intro and some fundamentals.

Going through books after this one was a breeze

I wish the README had a solid “what cool things you can do with this” right at the top.

In this day and age when programming is so accessible, why not have a more tempting pitch than just book titles categorized by difficulty.

I'll give you the TL;DR:

With CUDA, you can make Nvidia GPUs go brrrr.

Oh. And thereby, incidentally conquer the compute world.

First one I clicked on is 404: Programming Massively Parallel Processors: A Hands-on Approach (3rd Edition) https://www.cambridge.org/core/books/programming-in-parallel...

the newest is 4th ed i think

A fifth edition has been out recently: https://shop.elsevier.com/books/programming-massively-parall...

I started learning about GPU and CUDA from this book recently, and I agree the writing is confusing, and code examples have errors. However, it is still a nice reference about many types of algorithms for heterogeneous memory devices, it helped me understand better some patterns for CPUs.

Probably worth noting that writing performant kernels for modern Nvidia hardware looks almost nothing like what the books from 2012 are going to teach you. You can read them for fun if you'd like but they're basically irrelevant.

Does anyone know of any good resources for the newer paradigms like cuTile?

In an age when your company mandates you to raise your productivity right now with hundreds of percentage points using LLMs, how do you find an excuse to sit down and read a book?

It feels like a dirty secret, doesn't it?

Yeah, corps don't want you to know how to code, they want you to be a prompter...

Sometimes I squeeze in an hour or so a day to read. Living on the edge, looking for the next dopamine hit.

You guys have enough attention span left to read?

Don't you read while your agents are doing all the work for you? /s

Or make your agents do the reading for you!

Anthropunk

not on company time ?

As always, on private time, if available, otherwise wait when LLM connection breaks down.

Any good MOOCs on Parallel programming/NVIDIA?

https://www.youtube.com/playlist?list=PLzn6LN6WhlN06hIOA_ge6...

https://ppc.cs.aalto.fi/

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.

The story of Flash Attention is the best manifestation of power and difficulty of GPU programming. This page gives a nice overview of it https://aiwiki.ai/wiki/flash_attention

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.

Reply on news.ycombinator.com