The languages used here
The listings in this book are in four languages, and you do not need to be fluent in any of them to follow the argument. You do need to know which one you are looking at and what its notation is doing, because each one exists to express a different level of the same machine.
Most are CUDA C++. This is ordinary C++ with a small number of additions that mark which code runs where and how it is launched. A function annotated __global__ is a kernel: it runs on the GPU and is called from the host. The triple-chevron syntax on the call, kernel<<<blocks, threads>>>(args), is not standard C++ at all; it is the launch configuration, saying how many thread blocks to start and how many threads each contains. Inside the kernel, the built-in variables threadIdx, blockIdx, and blockDim tell each thread which one of those threads it is. That is enough to read the listing below, which is the whole shape of a CUDA program in miniature.
__global__ void scale(float* out, const float* in, float k, int n)
{
int i = threadIdx.x + blockDim.x * blockIdx.x;
if (i < n) out[i] = k * in[i];
}
scale<<<blocks, 256>>>(devOut, devIn, 2.0f, n);Two things in that listing are worth pausing on even before the C++ syntax. The first is that the kernel body contains no loop. The loop is the launch: rather than iterate over n elements, the program starts enough threads that each handles one element, and the index arithmetic on the first line is how a thread works out which one is its own. The second is the bounds check. The number of threads started is almost never exactly n, because threads come in fixed-size groups, so the last group runs past the end of the data and those threads must be told to do nothing.
C++, for Python readers#
If Python is your working language, five things in C++ will account for most of the friction.
Types are written down and fixed. float is a 32-bit number, int a 32-bit integer, and the compiler needs to know both before the program runs. There is no runtime type to inspect. This is not bureaucracy: knowing the exact size of every value is what lets the compiler decide how many of them fit in a register or a cache line, which is a question the later chapters ask constantly.
An asterisk marks a pointer. In float* out, the parameter is not an array of floats but the address of one. The kernel is handed a location in memory, not the data itself, and nothing about the type says how many elements live there or which memory it belongs to. That is why kernel signatures almost always carry a separate length argument, and it is the root of an entire class of bugs that Compute Sanitizer exists to catch. Pointers are also the reason this book talks so much about where memory lives: the address alone does not tell you, and reading from the wrong kind is a performance cliff rather than an error.
const is a promise. Writing const float* in says the code will not write through that pointer. The value of saying so is that the compiler may then keep the data in faster read-only paths, and reorder reads around it, in ways it cannot justify when a write might happen.
An ampersand marks a reference. Where a pointer is an address you can do arithmetic on, a reference is another name for an existing value, and passing one avoids copying a large object. In host code you will see both; in device code pointers dominate.
Angle brackets carry compile-time values. This is the one with the biggest payoff. A C++ template parameterizes code over values the compiler knows before the program runs, so a tile size written as Tile<128, 64> is not an argument passed at runtime but part of the type itself. The compiler can then unroll loops against it, allocate exactly the right registers, and discard branches that cannot be taken. That is the entire reason the CUTLASS family in CUTLASS, CuTe, and CUDA Tile is written the way it is, and why its error messages are so long: the shapes are in the types.
Python, and two kinds of output#
The second language is Python, in two dialects that look like Python and compile to GPU code. Triton kernels are Python functions under a @triton.jit decorator, and AWS’s NKI in Other hardware stacks follows the same idea on Trainium under @nki.jit. When a listing carries one of those decorators, the body is not being interpreted line by line; it is being compiled, and the Python you are reading is a description of a kernel rather than a program that runs.
The important difference is not the syntax but the unit of work. CUDA C++ asks you to write the code for one thread and then say how many to start. Triton asks you to write the code for one block of data, and the compiler decides how threads divide it up. The index arithmetic in the CUDA listing above simply does not appear, because working out which element belongs to which thread is no longer your job. That shift is the subject of the whole of Programming models and profiling, so it is enough here to recognize which of the two styles a listing is written in.
The last two languages are not written by anyone. PTX is the virtual instruction set NVIDIA compilers emit, and SASS is the actual machine code for one architecture. They appear in this book as evidence rather than as source: when a claim about what the hardware does needs proof, the compiler’s output is where the proof lives, because it shows what was actually issued rather than what the source appears to ask for. Compilation and machine code explains how one becomes the other and how to read either.
You can follow every argument in this book without running anything. If you do want to compile the CUDA listings, they go in a file ending .cu and are built with nvcc, which splits that single file into its host and device halves, compiles each, and links them back together. What you need installed is the CUDA Toolkit and an NVIDIA GPU to run the result on.