Week 8 · Lesson 1
How a GPU runs code
A desktop CPU has a few dozen hardware threads, each fast on its own. A GPU has thousands of simpler ones, and it's fast only when all of them have work. This lesson is the model you need to write kernels for it: blocks, threads, and the memories they share.
Many threads, grouped into blocks
A GPU program, a kernel, is one function that every thread runs. Threads come in blocks of up to 1024, and a launch starts a grid of blocks. Each thread knows two numbers: its block, blockIdx.x, and its position inside the block, threadIdx.x. Those numbers are how it picks its share of the work. A kernel that doubles 96 floats as 3 blocks of 32 threads has thread t of block b handle element 32b + t.
The hardware runs blocks on streaming multiprocessors (SMs). The RTX 4060 Ti this course was tested on has 34 of them. A block stays on one SM from start to finish, and the SM runs its threads 32 at a time, in groups called warps. Blocks run in whatever order the hardware likes, many at once, and a block can never wait for another block.
Three kinds of memory
| memory | who sees it | size and speed | in .uir |
|---|---|---|---|
| global | every thread of every block, and the host copies data in and out | gigabytes, hundreds of cycles away | buffers |
| shared | the threads of one block | tens of kilobytes per block, almost as fast as registers | an alloc used inside one global loop |
| registers | one thread | a few hundred bytes per thread, the fastest | values (the syllabus calls this ALU) |
Shared memory is what makes a GPU kernel fast. A block loads a tile of its inputs from global memory once, into shared memory, and then all its threads read the tile many times. That's week 7's tiling again, with the tile shared by a block of threads instead of kept by one.
Waiting for each other
The threads of a block can wait for each other with a barrier, __syncthreads(): no thread passes it until every thread of the block has reached it. That's how a block loads a tile and then uses it. Every thread writes its part, the barrier makes sure all parts are written, and then every thread may read any of them. Blocks have no such barrier, so two blocks must never depend on each other.
Why the bits still match
The grader compares your GPU's results with the interpreter, bit for bit, as it has every week. Each thread computes its outputs with the same operations in the same order as the CPU. Only two GPU habits would change bits, and the build rules out both: fused multiply-add, turned off with -fmad=false, and faster but less exact division, which -prec-div=true (already the default) keeps off. A real CUDA library wouldn't do that, and it would run a bit faster.
How fast?
The 4060 Ti does about 22 trillion float operations a second at its peak, and moves 288 GB a second from its global memory. Compare a CPU core: tens of billions. The catch is that all those threads need work and data at once. On a 512×512 GEMM, the tests measured:
| kernel (512×512 GEMM, RTX 4060 Ti) | time | speed |
|---|---|---|
| one thread per output, straight from global memory | 0.34 ms | 0.8 TFLOP/s |
| a block per row strip, a shared accumulator | 0.39 ms | 0.7 TFLOP/s |
| a 16×16 tile per block, loaded into shared memory | 0.17 ms | 1.6 TFLOP/s |
The row-strip kernel is slower than the naive one even though it uses shared memory: each of its 512 blocks has only a few threads, so the SMs have too few warps to keep busy while they wait for memory. The tiled kernel is the one the speed test uses as its reference. All three are far from 22 TFLOP/s. Getting closer takes the tricks of later weeks and of libraries like cuBLAS, and none of them can keep these exact bits.
You don't need a GPU for this week. Lesson 8.5 shows the free and cheap ways to get one, and the grader has an emulator that runs your CUDA on the CPU.
Check yourself
A kernel runs as blocks of 32 threads, and thread t of block b handles element 32b + t. Which element does thread 5 of block 2 handle?
Block 2 starts at element 64.
32·2 + 5 = 69. Blocks 0 and 1 cover elements 0 to 63.
Which threads can read a block's shared memory?
Shared memory lives on the SM running the block.
Each block gets its own copy of shared memory, on its SM, for as long as it runs. Other blocks can't see it. Data that must go from one block to another goes through global memory, in a later kernel.
Block 3 needs a value block 7 computes, in the same kernel. What's the problem?
What does __syncthreads() wait for?
A barrier waits for the threads of one block. Nothing in a kernel waits for another block, so block 7's results are only safe to read in a later kernel.
Why does the course build CUDA with -fmad=false?
How many roundings does a·b + c have?
a·b + c as two operations rounds the product, then the sum. A fused one rounds only at the end. That's more accurate, but different, and the grader wants the interpreter's exact bits.