Simon Boehm writes technical articles on GPU and CPU performance, distributed training, and programming practice.
Commenters repeatedly recommend his CUDA matmul article as one of the best write-ups on the gap between naive and tuned kernels.
This is a really good post, and a nice successor to the earlier one that I've pointed people at before: https://siboehm.com/articles/22/CUDA-MMM (this is linked in the post, of course). That one is a good introduction but I found that it doesn't really explain a lot of the choices it makes beyond the basic ones (coalescing=good, bank conflicts=bad). This is makes it a good introduction but this goes a lot further into the details that are important when you try to attack cuBLAS performance for faster kernels on more recent hardware.
In fact I recently had someone go through almost the same steps here, although they were using Ampere GPUs rather than Turing. As the post mentions, the microarchitecture is almost the same, especially when talking about the tensor cores (Hopper has significant changes). One difference is that it supports asynchronous global to shared memory transfers, which is quite useful. And you can get these for quite cheap: a RTX 3060 is a good card to test on. Vast.ai will rent one to you for a fraction of the price the author paid, and you can run Nsight Compute on them if you pick the VM instances (which costs several times more–the Docker instances are actually like a tenth of the prices the author paid).
I don't really have anything to point out about the content here but I will add that you will want to pipeline shared memory loads (to registers) in addition to those coming global memory. If you do this it's very possible to beat cuBLAS, which…
I know a trick, although it relies on the closed source cuBLAS library. The matrix multiplication in non-batched forward passes is really matrix-vector multiplication, but you can shoehorn that into matrix-matrix multiplication by setting the right dimension to 1. cuBLAS uses column major ordering to match Fortran instead of the row major ordering of C/C++, although you can swap transa and transb, m and n, a and b, lda and ldb; and if you are using cublasGemmEx() to use tensor cores, atype and btype. This relies on the property of mathematics where A * B = C and B’ * A’ = C’ for matrices. You can experiment with this on a CPU using a BLAS library like OpenBLAS or the Intel MKL before trying to do this on a GPU.
That said, to do matrix vector multiplication (GEMV), you just need to compute dot products on all of the rows of the matrix with the vector to get your output vector. You could probably just handle each dot product in power of 2 chunks on the GPU to compute partial sums. When you get to the last chunk that is not a power of 2, just have the threads that go pass the end of the dot product contribute 0 to the partial sums. Then you would have the threads handle the horizontal sum of the partial sums to finish the computation of an entry in the output vector.
As for matrix-matrix multiplication (GEMM), that is much harder to do in a performant manner. If you do it by treating the multiplication as a series of matrix-vector multiplications, you will be bandwidth…
Many of the optimizations here rely heavily on the size of matrix and it's relationship to hardware specific details, like LDS size, how they're banked and register count.
It's probably not surprising that you can grind a decent improvement over a general solution, and many of the improvements shown here will need to be re-balanced, or even simply not work, for kernels working on different matrix layouts. Similarly for trying to work on different hardware - even in the same architecture and generation these sort of details are often changing.
And all that required going down to the ISA level, which is a lot less easy (certainly less documented) for Nvidia - for example the "inspiration" post linked [0] on CUDA didn't beat cuBLAS also didn't try modifying the SASS directly, so there might be similar level gains unrealized there.
[0] https://siboehm.com/articles/22/CUDA-MMM
Modal's GPU glossary is a good overview about how GPUs work [0]. Karpathy's LLM overview is a good high level overview on LLMs [1]. 3b1b's video (and subsequent videos) on transformers was excellent at helping me understand the math at a high level [2]. This matrix multiplication optimization worklog helped me understand writing better CUDA (not for beginner intro though) [3].
During this process I also asked ChatGPT a lot of questions.
I'm definitely open to suggestions about "how to learn" with all the new tools we have. I felt this has not been straightforward to figure out.
[0] https://modal.com/gpu-glossary
[1] https://www.youtube.com/watch?v=7xTGNNLPyMI
[2] https://www.youtube.com/watch?v=wjZofJX0v4M
[3] https://siboehm.com/articles/22/CUDA-MMM
The code in this article is incorrect. The CUDA kernel is never called: https://github.com/RijulTP/GPUToolkit/blob/f17fec12e008d0d373e6e8978e6d26f312e2f4ce/mandelbrot/mandelGPU.py#L47
I'd also like to point out that 90 % of the time spent to "compute" the Mandelbrot set with the JIT-compiled code is spent on compiling the function, not on computation.
If you actually want to learn something about CUDA, implementing matrix multiplication is a great exercise. Here are two tutorials:
https://cnugteren.github.io/tutorial/pages/page1.html
https://siboehm.com/articles/22/CUDA-MMM
Have you tried to write a kernel for basic matrix multiplication? Because I have and I can assure you it is very hard to get 50% of maximum FLOPs, let alone 90%. It is nothing like CPUs where you write a * b in C and get 99% of the performance by the compiler.
Here is an example of how hard it is: https://siboehm.com/articles/22/CUDA-MMM
And this is just basic matrix mult. If you add activation functions it will slow down even more. There is nothing easy about GPU programming, if you care about performance. CUDA gives you all that optimization on a plate.