Chapter 2: First Light - Hello, CUDA
The machine had hummed in the corner of the classroom for a week, and for a week
nobody had been allowed to touch it. Guru-ji had spent the days at the chalkboard
instead: the post office with its thousand patient clerks, the rowing crews of
thirty-two, the wide river of memory. Arjun had endured all of it the way one
endures a long queue at a sweet shop, standing on one foot, eyes fixed on the
counter.
This morning Guru-ji finally pulled the wooden stool up to the keyboard. "Today,"
he said, "the machine speaks."
Arjun was at his elbow before the sentence ended. "Can we make it do a million
things at once? You said a thousand hands. I want to see the hands."
"You will see eight of them," said Guru-ji, "and then a million. But tell me
first - when you visit a neighbor's house, what do you do at the gate?"
Arjun blinked. "You call out. So they know you've come."
"So they know you've come," Guru-ji agreed. "The card inside this machine is a
neighbor's house. It has its own rooms, its own shelves, its own workers. Our
first program will do nothing but stand at the gate and ask every worker to call
back. When you hear eight voices answer, beta, you will know the gate is open."
Kabir leaned in from the other side. "And the million things?"
"Before lunch," said Guru-ji, and reached for the switch.
In Chapter 1 (A Thousand Hands) we toured the GPU from the outside: thousands of
small cores organized into streaming multiprocessors, warps of 32 threads pulling
in lockstep like a rowing crew, and a memory system whose bandwidth is the real
currency of everything we will build. It was all talk. This chapter is where you
type.
We will write two complete programs. The first makes the GPU print a greeting.
Trivial, except that it forces us through every rite of passage at once: the
split between the two processors in your machine, the strange triple-bracket
launch syntax, the compiler that serves two masters, and the discipline of
checking every call for errors. The second program adds two arrays of a million
floating-point numbers each, which sounds equally trivial until you realize that
a million additions happening at once is precisely the shape of work a neural
network is made of. The kernel you write today is a direct ancestor of the ones
that will train our model in train_gpt.cu and, in Chapter 21, write its first
line of Shakespeare.
By the end you will be able to allocate memory on the GPU, move data to it and
back, run your own function across thousands of threads, and - this matters most
- prove the answer is right. That last habit starts today and never stops.
Two computers in one box
The first thing to understand about CUDA is that you are never programming one
computer. You are programming two, and they are unequal partners living in the
same case.
One of them is the machine you already know: the CPU, its handful of fast cores,
and the system RAM your malloc hands out. CUDA calls this side the host.
The other is the graphics card: its own processor with thousands of small cores,
and (this is the part newcomers miss) its own separate memory, soldered onto
the card itself. CUDA calls this side the device. Device memory is fast in a
way that makes system RAM look sleepy: on a typical modern GPU it moves around
1 TB/s, against roughly 50 GB/s for the CPU and its RAM.
Between the two sides runs a bridge: the PCIe bus. And the bridge is slow. On a
common PCIe 4.0 machine it carries about 32 GB/s, a thirtieth of what the GPU's
own memory can do. Every byte the GPU works on must first cross this bridge, and
every result you want to see must cross back.
HOST (the CPU side) DEVICE (the GPU side)
+------------------------+ +--------------------------+
| CPU | | GPU |
| a few fast cores | PCIe | thousands of small |
| | bridge | cores |
| system RAM |<========>| device memory |
| ~50 GB/s to the CPU | ~32 GB/s | ~1 TB/s to the GPU |
+------------------------+ +--------------------------+
your malloc lives here cudaMalloc lives here
Guru-ji's image of the neighbor's house is worth keeping. The GPU is not an
extra room in your house; it is the house next door, with its own kitchen and
its own pantry. If you want the neighbors to cook for you, you must carry the
ingredients over, knock, wait, and carry the finished dish back. The carrying is
expensive, so the whole art of CUDA programming - most of this book, really -
is arranging for a great deal of cooking per trip.
This split shapes the structure of every CUDA program you will ever write. The
host is the manager: it allocates, copies, launches work, and collects results.
The device is the workforce: it does the actual arithmetic, in bulk, on data
that already lives on its side of the bridge. Your source file contains code for
both, side by side, and part of learning CUDA is keeping straight which lines
run where.
One consequence deserves its own paragraph, because it produces the most
baffling beginner crashes: a pointer returned by cudaMalloc points into device
memory, and the host cannot dereference it. The number looks like any other
pointer. Print it and you see a plausible address. But d_a[0] on the host is a
trip to a house you cannot enter, and the program dies, or worse, doesn't die
and reads garbage. We will keep device pointers visually distinct with a d_
prefix throughout the book, a small convention that has saved me more debugging
hours than any tool.
Before CUDA, a beautiful hack
Making a graphics card do general math did not begin gracefully. Before CUDA arrived in 2007, researchers who wanted the GPU's arithmetic had to disguise their problem as a graphics problem: they packed their numbers into textures (images), wrote their math as a "shader" (a tiny program meant to colour pixels), and read the answers back out as rendered images. It worked, barely, and it was miserable.
CUDA's quiet revolution was to let you stop pretending - to write plain C that says "run this function on the GPU" and mean exactly that. The
__global__keyword you meet in this chapter is that whole revolution, compressed to one word.
A kernel is a function that runs N times at once
On the CPU, calling a function runs it once. The GPU's core idea is different
enough that it deserves a slow sentence: a kernel is a function that you
write once and launch as many thousands of copies of as you like, all running
at the same time, each copy knowing only one thing the others don't: its own
index.
Think of Guru-ji handing out an errand, not to one student but to the whole
village school at once. He does not write a different note for each child. He
writes one note - "go to the house whose number you were given, and sweep its
doorstep" - and hands every child a number. Same instructions, different house.
The note is the kernel; the number is the thread index; the sweeping is your
arithmetic.
In CUDA, each running copy of the kernel is called a thread. Threads come
grouped into thread blocks (teams that will later learn to cooperate), and
all the blocks of one launch together form the grid. Chapter 3 (An Army in
Formation) is devoted to this hierarchy and why it exists; today we only need
the vocabulary and the two variables CUDA gives every thread so it can find its
number: threadIdx.x, its position within its block, and blockIdx.x, which
block it belongs to.
Marking a function as a kernel takes one keyword. Write __global__ in front
of it, and the compiler knows: this function's body runs on the device, but the
call comes from the host. A kernel always returns void: with thousands of
copies running at once, there is no sensible single return value; results come
back through memory instead.
Launching a kernel uses syntax you will not have seen in any C book, because it
is not C - it is CUDA's one big extension to the language:
my_kernel<<<blocks, threads_per_block>>>(arguments);
The triple angle brackets carry the launch configuration: how many blocks
to create, and how many threads in each block. Launch with <<<2, 4>>> and you
get 2 blocks of 4 threads: 8 copies of the function. Launch with
<<<3907, 256>>> and you get slightly over a million copies. The arguments in
the ordinary parentheses are passed to every copy identically, just like the
wording on Guru-ji's note.
That is the entire conceptual core of CUDA: one function, many indexed copies,
launched over data that already lives on the device. Everything else in this
book - tiling, reductions, attention - is arrangement and choreography on top
of this one idea.
The two most famous words in programming
Every programmer's first program prints "Hello, World". The phrase was made famous by Brian Kernighan, who used it in a 1972 Bell Labs tutorial for the B language and then enshrined it in 1978 in The C Programming Language, the book so universally known that programmers just call it "K&R" after Kernighan and its co-author Dennis Ritchie, C's inventor.
Half a century later, running those two words in a new language is still the rite of passage that says: the toolchain works, and I am really here. This chapter's version happens on ten thousand cores at once.
Hello from the device
Here is the whole first program (full program in code/ch02/hello.cu):
/*
hello.cu - the first program that runs on the GPU (Chapter 2).
Launches a kernel on 2 blocks of 4 threads each; all 8 threads print who
they are. Build with: make hello
*/
#include <stdio.h>
#include "../common/common.h"
// __global__ marks a function that runs on the device (GPU) but is
// launched from the host (CPU). Every kernel returns void.
__global__ void hello_kernel(void) {
printf("Hello from block %d, thread %d\n", blockIdx.x, threadIdx.x);
}
int main(void) {
hello_kernel<<<2, 4>>>(); // 2 blocks x 4 threads = 8 hellos
CUDA_CHECK(cudaGetLastError()); // did the launch itself fail?
CUDA_CHECK(cudaDeviceSynchronize()); // wait for the GPU, flush its printf
return 0;
}
Walk through it from the top. The include pulls in common.h, the one shared
header of this book's code; we will open it up in a moment to see what
CUDA_CHECK actually is. Then comes the kernel: __global__ void
hello_kernel(void). Its body is a single printf, and yes, that printf
runs on the graphics card. The device runtime collects the output of every
thread into a buffer on the GPU and delivers it to your terminal later. Notice
what gets printed: blockIdx.x and threadIdx.x, the two coordinates that
tell each thread who it is. The kernel takes no arguments and computes nothing;
its only job is to prove that many copies really ran.
In main, the launch hello_kernel<<<2, 4>>>() creates 2 blocks of 4 threads:
8 threads in all, each executing the same one-line body with its own
coordinates. Then two lines that look like ceremony and are anything but. A
kernel launch is asynchronous: the host does not wait for the GPU. The
launch line drops the work into the GPU's queue and returns immediately, in
microseconds, while the device gets on with it. If main simply returned at
that point, the program could exit before the GPU had printed a thing.
cudaDeviceSynchronize() is the host saying: block right here until the device
has finished everything I have asked of it. It is also the moment the device's
printf buffer gets flushed to your screen.
The cudaGetLastError() call just before it deserves its own mention. The
launch syntax has no return value you can inspect, so CUDA records launch
failures (a block size too large, no device present) in a hidden error slot
instead. Asking for the last error immediately after a launch is how you find
out the launch even happened. Silent launch failures are the source of the
classic "my kernel is infinitely fast" bug, which Arjun will personally
rediscover in Chapter 7 (Trust but Verify).
Build and run it, and eight voices answer:
$ make hello
$ ./hello
Hello from block 0, thread 0
Hello from block 0, thread 1
Hello from block 0, thread 2
Hello from block 0, thread 3
Hello from block 1, thread 0
Hello from block 1, thread 1
Hello from block 1, thread 2
Hello from block 1, thread 3
One caveat: the order is not guaranteed. Blocks run whenever the
hardware schedules them, so your block 1 may well speak before your block 0.
The tidy ordering above is common on an idle GPU but it is luck, not law, and
the sooner you stop expecting threads to run in order, the sooner half of your
future bugs evaporate. The gate, in any case, is open.
nvcc, briefly
You cannot build this with gcc or clang alone, because half the file is not
CPU code. NVIDIA's compiler driver nvcc splits your .cu file in two: the
host code goes to your regular system compiler, and the device code (everything
__global__, plus the launch syntax) is compiled for the GPU. The two halves
are then stitched back into one ordinary executable; the device code rides
along inside it in a container called a fat binary, which can hold compiled
code for several GPU generations at once so one executable runs on many cards.
The Makefile in code/ch02/ invokes it for you:
NVCC ?= nvcc
NVCCFLAGS ?= -O3 -arch=native -lm
-O3 is the usual optimization level, and -arch=native tells nvcc to compile
for whatever GPU is in the machine doing the compiling, the right default for
a book where you build and run on the same box. We explicitly leave fast-math
off: it trades exact IEEE floating-point behavior for speed, and this book's
verification habit depends on results we can compare bit for bit against the
CPU. That is genuinely all you need to know about the toolchain for now;
Appendix B covers installation and the cloud-GPU route from zero.
The habit of checking everything
Nearly every function in the CUDA runtime API returns an error code of type
cudaError_t. Nearly every beginner ignores it. The result is a special kind
of misery: the allocation that silently failed three hundred lines ago finally
detonates inside an innocent kernel, and you spend an evening debugging the
wrong function.
The fix costs one macro, defined in code/common/common.h and used on every
CUDA call in this book without exception:
static inline void cuda_check(cudaError_t err, const char* file, int line) {
if (err != cudaSuccess) {
fprintf(stderr, "[CUDA ERROR] %s:%d: %s\n",
file, line, cudaGetErrorString(err));
exit(1);
}
}
#define CUDA_CHECK(call) cuda_check((call), __FILE__, __LINE__)
Mechanically, this is plain C. The macro wraps any expression that yields a
cudaError_t and hands the result to cuda_check, along with the source file
and line number that the preprocessor fills in via __FILE__ and __LINE__.
If the call succeeded, nothing happens and the program flows on. If it failed,
you get the exact file, the exact line, and a human-readable message from
cudaGetErrorString (something like out of memory or invalid device
ordinal), and the program stops immediately, at the true scene of the crime
rather than three hundred lines downstream.
Exiting on the spot may look drastic. For a long-running service it would be;
for the programs in this book it is exactly right, because a failed CUDA call
means every result after it is fiction, and the kindest thing a program can do
with fiction is refuse to print it. You will meet the rest of common.h (the
GPU timer, the random number generator for weight initialization) as the
chapters need them, and Chapter 7 gives the whole error model a proper tour.
For now, adopt the rule as written: every call wrapped, no exceptions, from
day one. It is the cheapest insurance in GPU programming.
Memory on the other side of the bridge
Printing proved the workers exist. To make them compute, we must first give
them something to compute on, and that means learning the three calls that
manage the neighbor's pantry: cudaMalloc, cudaMemcpy, and cudaFree.
cudaMalloc is malloc for device memory. Its signature trips people up the
first time:
float* d_a;
float* d_b;
float* d_out;
CUDA_CHECK(cudaMalloc(&d_a, size));
Where malloc returns the new pointer, cudaMalloc returns an error code
(everything in CUDA returns an error code), so the pointer comes back through
an output parameter. You pass the address of your pointer, and the runtime
fills it in with the location of size fresh bytes of device memory. Forget
the & and the compiler will usually save you; cast your way past the
compiler and nothing will.
cudaMemcpy is the truck that crosses the bridge:
CUDA_CHECK(cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice));
Destination first, source second, like memcpy. The final argument states the
direction of travel: cudaMemcpyHostToDevice carries ingredients over,
cudaMemcpyDeviceToHost brings the dish back. The direction must match the
pointers you pass: the runtime cannot reliably guess which side of the bridge
an address lives on, and lying to it is undefined behavior of the most
creative kind.
Three properties of cudaMemcpy matter enough to state now. First, it is
synchronous: when the call returns, the copy is complete. Second, it quietly
waits its turn: if a kernel is still running on data involved in the copy,
the copy begins only after the kernel finishes. This is why the vector-add
program below needs no explicit cudaDeviceSynchronize: the copy-back doubles
as one. Third, the bridge crossing is slow enough that
you should picture it every time you type the word cudaMemcpy. Our two input
arrays are 4,000,000 bytes each; at about 32 GB/s the 8 MB outbound trip costs
roughly a quarter of a millisecond, which will turn out to be far longer than
the million additions themselves take. Chapter 4 (The Memory Hierarchy) makes
a science of this; for now it is enough to feel the asymmetry.
cudaFree releases device memory, and it gets wrapped in CUDA_CHECK like
everything else. Leaking device memory is easier than leaking host memory,
because no tool you already use watches for it, and a GPU that slowly fills up
across runs of a buggy program is a confusing thing to debug.
A million additions at once
Now the real program. We want out[i] = a[i] + b[i] for a million values of
i, and we want every addition to happen on the device. On the CPU this is a
loop. On the GPU, the loop disappears: we launch a million threads and each
one performs a single addition. Get this one move right and you have the engine
of the entire book: the layernorms, attention scores, and gradients that will
train train_gpt.cu are all, underneath, vast piles of independent arithmetic
exactly like this, each element waiting for a thread of its own. Here is the
kernel (full program in
code/ch02/vector_add.cu):
// Each thread computes exactly one element: out[i] = a[i] + b[i].
__global__ void vector_add_kernel(float* out, const float* a, const float* b,
int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) { // the grid overshoots; stay in bounds
out[i] = a[i] + b[i];
}
}
Four lines of body, and two of them are the most important lines in this
chapter. Take the first one slowly, because every kernel you will ever write
begins with some variation of it:
int i = blockIdx.x * blockDim.x + threadIdx.x;
A single block can hold at most 1,024 threads, so a million threads must be
spread across many blocks, and no thread is directly told its global position.
Each thread knows three things: which block it is in (blockIdx.x), how big
every block is (blockDim.x, a built-in that the launch configuration fills
in; 256 for us), and its seat within its own block (threadIdx.x). The
formula converts those three into a unique global index, exactly like a street
address. If every lane in the village has 256 houses, then the house at
position 5 in lane 2 is house number 2 * 256 + 5 = 517 counting from the
village gate: lanes 0 and 1 contribute 512 houses before ours, and we stand 5
doors into lane 2.
The second important line is the bounds check, if (i < n), and the reason it
must exist is a small arithmetic story worth telling properly.
Covering the array with blocks
We chose 256 threads per block: the book's default, a multiple of the
32-thread warp and a size that behaves well on every modern GPU. How many
blocks do we need to cover n = 1,000,000 elements? Dividing gives
1,000,000 / 256 = 3,906.25, and here is the trap: integer division in C
truncates, so n / block_size yields 3,906 blocks, which is
3,906 × 256 = 999,936 threads. The last 64 elements of the array would never
be touched, and nothing would warn you. Meera's kind of bug: the answer is
almost right.
The fix is to round the division up, and the idiom for rounding up in
integer arithmetic is worth committing to memory, because it appears in nearly
every CUDA program ever written:
int block_size = 256;
int grid_size = (n + block_size - 1) / block_size; // 3907 blocks
Adding block_size - 1 before dividing nudges any nonzero remainder over the
edge into one extra block, while leaving exact multiples alone. Check both
cases with real numbers. For n = 1,000,000: (1,000,000 + 255) / 256 =
1,000,255 / 256, which truncates to 3,907 - one block more than the naive
division, and 3,907 × 256 = 1,000,192 threads, enough to cover everything.
For a hypothetical n = 999,936, an exact multiple of 256: (999,936 + 255) /
256 = 1,000,191 / 256 truncates to 3,906, so no extra block is wasted. The
idiom is called ceiling division, and from here on the book uses it without
comment.
Rounding up creates an overshoot: 1,000,192 threads for 1,000,000 elements
means the last 192 threads compute an index i of 1,000,000 or more, pointing
past the end of every array. Without the bounds check, those threads would
write into memory that belongs to something else - perhaps harmlessly, perhaps
corrupting a neighboring array, perhaps crashing, depending on luck and
layout. With the check, they compute their index, find it too large, and
quietly do nothing. A few idle threads out of a million is a price too small
to measure.
array: n = 1,000,000 floats (indices 0 .. 999,999)
+--------+--------+--------+-- ... --+-----------+
| 256 | 256 | 256 | | 256 |
+--------+--------+--------+-- ... --+-----------+
block 0 block 1 block 2 block 3906
i = 0 256 512 999,936
..255 ..511 ..767 ..1,000,191
^ the last 192
threads: i >= n,
do nothing
Guru-ji wandered over while Meera was tracing the last block's indices on her
slate. "So some children are sent to houses that do not exist," he said.
"A hundred and ninety-two of them," said Meera. "They knock on air."
"And what should a child do, finding no house?"
"Nothing. Turn around, go home." She frowned at her slate. "It seems wasteful."
"It is the opposite - it is the bargain of the season," said Guru-ji. "To give
every real house a visitor, you must hire by the lane, not by the house. A few
children stroll for nothing, and in exchange no house is ever missed. Cheap.
Now, the dangerous question: what happens if a child forgets to check, and
sweeps the doorstep of a house that isn't there?"
Meera looked up from her slate, and did not need to answer.
The host side, top to bottom
The kernel is four lines; the host code around it is the ritual you will
perform, with variations, in every program for the rest of the book. Here is
the first half of main - allocate, fill, ship:
int main(void) {
int n = 1000000; // one million elements
size_t size = n * sizeof(float); // 4,000,000 bytes per array
// host memory: two inputs, the GPU result, and a CPU reference
float* a = (float*)MALLOC_CHECK(size);
float* b = (float*)MALLOC_CHECK(size);
float* out = (float*)MALLOC_CHECK(size);
float* ref = (float*)MALLOC_CHECK(size);
for (int i = 0; i < n; i++) {
a[i] = (float)i * 0.001f; // deterministic, easy to reason about
b[i] = (float)(n - i) * 0.001f;
}
// device memory: mirrors of a, b, and out on the other side of the bridge
float* d_a;
float* d_b;
float* d_out;
CUDA_CHECK(cudaMalloc(&d_a, size));
CUDA_CHECK(cudaMalloc(&d_b, size));
CUDA_CHECK(cudaMalloc(&d_out, size));
// host -> device: the inputs cross the bridge once
CUDA_CHECK(cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice));
We allocate four host arrays, not three: the two inputs, a buffer for the
GPU's answer, and ref for the CPU's answer, because we intend to compare
them. MALLOC_CHECK is common.h's wrapper around malloc that aborts with
a file and line number if the allocation fails: the same philosophy as
CUDA_CHECK, applied to the host.
The fill values are chosen to be boring on purpose. a[i] climbs from 0 in
steps of a thousandth; b[i] descends the same way; their sum should sit near
1,000 everywhere. Deterministic data means a failed check is reproducible: run
it again and the same element fails with the same numbers, which is the
difference between an evening of debugging and a week of it. Then come the
three cudaMalloc calls for the device mirrors and two cudaMemcpy calls to
carry the inputs across. Note what does not cross: d_out is allocated on
the device but nothing is copied into it, because the kernel will overwrite
every element. The result buffer is born on the far side of the bridge.
Next comes the second half, launch and retrieve:
// one thread per element; round the block count UP so no element is missed
int block_size = 256;
int grid_size = (n + block_size - 1) / block_size; // 3907 blocks
vector_add_kernel<<<grid_size, block_size>>>(d_out, d_a, d_b, n);
CUDA_CHECK(cudaGetLastError());
// device -> host: cudaMemcpy waits for the kernel to finish first
CUDA_CHECK(cudaMemcpy(out, d_out, size, cudaMemcpyDeviceToHost));
The launch hands device pointers and n to every one of the 1,000,192
threads; cudaGetLastError confirms the launch was accepted, exactly as in
hello.cu. One cudaMemcpy in the other direction then brings the answer
home. Recall the ordering guarantee from two sections back: the copy waits for
the kernel to finish before it starts reading d_out, so by the time this
line returns, out on the host holds the complete result. No explicit
synchronize is needed; the copy is the synchronize.
Cleanup closes the file: all seven allocations are freed, device side wrapped
in CUDA_CHECK, host side with plain free (see the last dozen lines of
code/ch02/vector_add.cu). Between the copy-back and the frees, though, sits
the most important part of the file, and it is Meera's.
Meera's rule: never trust an unchecked answer
It would be easy to print out[0], see a plausible number, and declare
victory. It would also be exactly how subtle GPU bugs survive: an indexing
mistake that scrambles the middle of the array leaves out[0] perfectly
correct. A million elements came back from a processor we cannot watch, and we
have looked at one of them.
So the program computes the entire answer a second time, on the CPU, the slow
and obvious way, and compares every single element:
// CPU reference: the same million additions, one at a time
for (int i = 0; i < n; i++) {
ref[i] = a[i] + b[i];
}
check_close(out, ref, n, 1e-5f, "vector_add");
check_close lives in common.h and is the last function every numeric
program in this book calls:
static inline int check_close(const float* a, const float* b, size_t n,
float atol, const char* name) {
for (size_t i = 0; i < n; i++) {
float diff = fabsf(a[i] - b[i]);
if (diff > atol || isnan(diff)) {
printf("FAIL %s: index %zu: %f vs %f (diff %f, atol %f)\n",
name, i, a[i], b[i], diff, atol);
return 0;
}
}
printf("OK %s\n", name);
return 1;
}
It scans both arrays and demands that every pair agree to within atol, the
absolute tolerance. On success it prints one line, OK vector_add, and on
failure it prints the first offending index with both values, which is usually
all the forensic evidence you need: a mismatch at index 999,936 points
straight at a last-block problem, and a mismatch at every odd index points at
an indexing formula. The isnan guard matters too: NaN, the floating-point
"not a number" that arithmetic on garbage produces, compares false with
everything, so diff > atol alone would wave NaNs through.
Why a tolerance at all, and why 1e-5f here? Floating-point addition rounds,
so two computations of the same sum can disagree in the last bits if they do
anything differently. In this program they don't: GPU and CPU perform the same
single IEEE addition on the same two floats, so the results should match
exactly, and the tolerance is pure margin. Starting with Chapter 6 (When
Threads Disagree), our kernels will add numbers in a different order than
the CPU loop does, and order changes rounding; the book's standard tolerance
becomes 1e-4 then, and Chapter 10 makes the whole business of fair float
comparison precise. Meera would want you to know the tolerance is a choice,
not a magic number.
Run it:
$ make vector_add
$ ./vector_add
OK vector_add
One line of output for a million verified additions. Underwhelming to look at,
and the foundation of everything: from this chapter to the last, no kernel in
this book is declared working until a CPU reference agrees with it. When we
assemble the full model, its layernorms and softmaxes and attention kernels will
each carry a check like this one, and the finite-difference gradient checker
of Chapter 17 is this same idea grown teeth. The habit costs a dozen lines per
program. Skipping it costs weekends.
The terminal printed OK vector_add and the classroom exhaled. Arjun,
naturally, was already reaching for the stopwatch he keeps in his head. "A
million additions. How long did it take? Can we race the CPU properly?" Guru-ji
closed the lid of his Shakespeare, which had sat beside the keyboard all
morning like a promise. "You built one lane of children today, then thirty-nine
hundred lanes, and every doorstep was swept and checked. Tomorrow we ask how
the lanes are drawn up - and after that, Arjun, you may bring your stopwatch."
What we built
Two small programs, and every habit that the rest of the book stands on. The
GPU is no longer an abstraction from Chapter 1's chalkboard; it is a second
computer you can allocate on, copy to, compute with, and, most importantly,
verify against.
code/ch02/hello.cu: a__global__kernel launched with<<<2, 4>>>,
printing from 8 device threads;cudaDeviceSynchronizeto wait for the GPU.code/ch02/vector_add.cu: one million float additions, one thread per
element;cudaMalloc,cudaMemcpyin both directions,cudaFree.- The global index formula
blockIdx.x * blockDim.x + threadIdx.xand the
ceiling-division idiom(n + block_size - 1) / block_sizefor sizing grids. - The bounds check
if (i < n)that keeps the overshooting threads harmless. CUDA_CHECKon every API call andcudaGetLastErrorafter every launch.- The verification habit: a CPU reference and
check_close, ending inOK.
Where we're headed: our million-element launch needed 3,907 blocks, and we
waved that number into existence with one line of arithmetic. But why must
threads come bundled into blocks at all? Who decides which block runs first,
and where? And how does the picture change when the data is a matrix rather
than a line? That is Chapter 3 (An Army in Formation - Threads, Blocks, and
Grids), where the thread hierarchy stops being vocabulary and becomes a tool.
Exercises
-
In
hello.cu, change the launch configuration from<<<2, 4>>>to
<<<4, 2>>>. Before running it, write down exactly what the 8 lines of
output will say. Then run it. Did the order match your prediction, and
should it? -
In
vector_add.cu, changento 1,000,001 and recomputegrid_sizeby
hand. How many blocks are launched, and how many threads do nothing? Then
delete theif (i < n)bounds check, rebuild, and run a few times. Explain
why the program might printOK vector_addanyway - and why that makes the
missing check more dangerous, not less. -
Modify
vector_add.cuto computeout[i] = a[i] * b[i] + a[i]instead of
the sum, updating both the kernel and the CPU reference. Witha[i]as
large as 1,000, is1e-5fstill a reasonable absolute tolerance? Try it
and reason about what you see. -
Write a new program,
fill_squares.cu, whose kernel takes onlyfloat* out
andint nand setsout[i] = (float)i * (float)i- no input arrays, so
nothing needs copying to the device. Verify against a CPU loop with
check_close. Reuse the Makefile pattern to add a target for it. -
For the curious (Arjun has already done this): peek ahead at
GpuTimerin
code/common/common.hand wrap the kernel launch intimer_startand
timer_stop, then wrap the two host-to-device copies the same way. Which
costs more, the million additions or the trip across the bridge? Chapter 4
explains what you just measured.