CUDA basics: the words and the picture

Read this first

This guide explains the words you will meet in the CUDA Playground, in the order you need them. You do not need a GPU, and you do not need to know C. Each section is short and ends with something to try. Plan about 15 minutes, then open the CUDA Playground and run the examples.

1. What CUDA is

CUDA is NVIDIA's way of writing programs that run on a graphics card (GPU). A CPU has a few powerful cores that work through jobs quickly, one after another. A GPU has thousands of small cores that do simple jobs at the same time. CUDA code is C++ with a few extra keywords. It helps most when the same small operation must be done on many items: pixels, numbers in an array, rows of a matrix.

2. Host and device

• Host is the CPU and its normal memory (RAM).
• Device is the GPU and its own memory. Data is copied between the two, and a kernel can only work on data that is already in device memory.

3. The six steps of a CUDA program

1. Prepare data on the hostThe CPU fills arrays in its own memory.
2. Allocate GPU memorycudaMalloc reserves space in the device memory.
3. Copy the inputs overcudaMemcpy with cudaMemcpyHostToDevice moves the data to the GPU.
4. Launch the kernelkernel<<<blocks, threadsPerBlock>>>(…) starts the threads.
5. Copy the results backcudaMemcpy with cudaMemcpyDeviceToHost brings the answer to the CPU.
6. Free GPU memorycudaFree releases what cudaMalloc reserved.

The playground does steps 1, 2, 3, 5 and 6 for you. Step 4, the kernel, is where most mistakes happen, so that is where you should spend your time.

4. Your first kernel: vector addition

On a CPU, a loop adds the elements one at a time. On a GPU the loop goes away: each thread adds one pair.

// CPU: one loop does every element
for (int i = 0; i < n; i++) c[i] = a[i] + b[i];

// GPU: the kernel does one element per thread
__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;   // which element is mine?
  if (i < n) c[i] = a[i] + b[i];                   // only if it exists
}

// Launch enough blocks of 256 threads to cover n elements
vecAdd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);

• __global__ marks a kernel: a function that is called from the CPU and runs on the GPU.
• The index line gives each thread its own element. Section 6 explains it.
• if (i < n) stops the spare threads from touching memory past the end. Section 8 explains it.

5. Thread, block, grid

When you launch a kernel you ask for a number of threads. They are arranged in two levels:

• Thread: one run of the kernel function. Every thread runs the same code on different data.
• Block: a group of threads, up to 1,024. Threads in one block can share fast memory and wait for each other. Blocks run independently and in any order.
• Grid: all the blocks of one launch.

Take a launch of 3 blocks with 8 threads each. That is 24 threads. Each block has threads 0 to 7, and each thread knows where it is through four built-in values:

6. Finding your element: the global index

Each thread needs one number that is unique across the whole grid. The formula is:

int i = blockIdx.x * blockDim.x + threadIdx.x;

Worked example: thread 5 of block 2, with 8 threads per block, gets 2 × 8 + 5 = 21. The blocks cover the elements in order:

Every element is handled by exactly one thread. Try it: click a thread in the playground and check its index against this formula.

7. The launch configuration: <<<blocks, threads>>>

The triple angle brackets tell CUDA how many threads to start:

vecAdd<<<3, 8>>>(a, b, c, n);   // 3 blocks of 8 threads each

The values in brackets are the launch configuration. The values in parentheses are the arguments that every thread receives.

In the playground, step 2, Launch Configuration, has the same three things:

• The two number boxes are blocks and threads per block. The header shows the total thread count and the warp count.
• Fit blocks to N works out how many blocks cover N elements. The formula is blocks = (n + threadsPerBlock − 1) / threadsPerBlock, which rounds up.
• The argument rows set up the arrays. "length" is how many elements each array has, "starts as" fills it with values, and int n is the number of real elements.

8. Spare threads and the bounds check

Blocks come in whole sizes. With 20 elements and 8 threads per block you need 3 blocks, which is 24 threads. So 4 threads have no element. Without a check, those threads read or write past the end of the array, and that is undefined behaviour: the result can be wrong or the program can crash. if (i < n) makes the spare threads do nothing. The playground's "Missing bounds check" example shows the damage.

9. Warps and lanes

The GPU does not run threads one at a time. It groups 32 consecutive threads of a block into a warp, and all 32 run the same instruction at the same moment. Each position inside a warp is a lane.

A block of 8 threads still gets a whole warp, but only 8 of its 32 lanes do work. The other 24 sit idle. That is why block sizes should be multiples of 32, such as 64, 128 or 256. The playground shows one row per warp, so you can see the idle lanes.

10. Warp divergence

All threads of a warp run the same instructions. If some threads take an if and others take the else, the warp runs the first side with the other threads switched off, and then runs the second side. Both sides cost time.

When a whole warp takes the same side, there is no cost. The "Warp divergence" example shows the slow version. "No divergence" shows the same work arranged so that each warp agrees.

11. Race condition

Threads run in parallel, with no fixed order. If two threads write the same element, one write can overwrite the other. The final value then changes from run to run. Two fixes: atomicAdd, which adds safely even when threads collide, or giving each thread its own output element. The playground marks elements that were written by more than one thread.

12. Streaming multiprocessor (SM)

An SM is the hardware unit that runs blocks. A real GPU has dozens to over a hundred of them, and the blocks of your launch are spread across them. You do not choose which SM runs a block, so blocks must not depend on each other's order.

13. Reading the playground

• 1 Kernel: the code. The header explains the example, and the numbered comments explain each step.
• 2 Launch Configuration: blocks, threads per block, and the arrays and n that the kernel receives.
• 3 Threads: every thread as a square, one row per warp. Click a square to see its index, its element and what it wrote. Colours: wrote, did nothing, error (an out-of-bounds access), wrote where another thread also wrote (a race), and divergent warp.
• 4 Memory: the arrays before and after the run, with the elements each thread wrote.

Words in one place

Every term used on this guide and in the playground, in alphabetical order.

atomicAdd

A function that adds to a value safely when several threads add to the same element at once.

Block

A group of up to 1,024 threads. Its threads can share fast memory and wait for each other.

blockDim, blockIdx, threadIdx, gridDim

Built-in values that tell a thread where it is: the block size, its block's number, its own number in the block, and the number of blocks.

Bounds check

An if that stops a thread from using an index past the end of an array, usually if (i < n).

cudaMalloc, cudaMemcpy, cudaFree

Functions that reserve GPU memory, copy data between host and device, and release the memory again.

Device

The GPU and its memory.

Global index

blockIdx.x * blockDim.x + threadIdx.x: one unique number per thread across the whole grid.

Grid

All the blocks of one kernel launch.

Host

The CPU and its memory.

Kernel

A function marked __global__. It is called from the CPU and runs once on every thread of the launch.

Lane

One of the 32 positions in a warp. A lane is idle when its thread does not exist or has no work.

Launch configuration

The numbers in kernel<<<blocks, threads>>>: how many blocks, and how many threads in each.

Race condition

Two or more threads writing the same element without coordination. The result can change from run to run.

Shared memory

Fast memory that all threads of one block can use. Not yet in the playground.

Streaming multiprocessor (SM)

The hardware unit that runs blocks. The GPU spreads a launch's blocks across its SMs.

Thread

One run of the kernel function, with its own index.

Warp

A group of 32 consecutive threads of a block that run the same instruction together.

Warp divergence

When threads of one warp take different sides of an if. The warp runs each side in turn, which costs time.

Questions