Introduction to CUDA

A GPU was first built to draw graphics. To draw a picture, it has to calculate the color of millions of pixels, and it can do many of these small calculations at the same time. CUDA is a platform created by NVIDIA that lets us use this power for our own programs, not only for graphics. It works with NVIDIA GPUs. In this content, we will see the main idea behind it without installing anything: how a big job is split into many tiny pieces.

A CPU has a small number of powerful cores. A GPU has thousands of simpler ones. Imagine we need to paint every pixel of an image. The CPU paints a few pixels at a time, but very quickly. The GPU paints many pixels at once, even though each of its cores is slower. Run both of them, then try the smallest size.

When the job is small, the CPU wins. Before the GPU can start, the data usually has to be copied to its own memory, and the work has to be started. With only a few pixels, this setup takes longer than the work itself. As the job grows, the GPU leaves the CPU far behind. That is why GPUs are great for work that can be split into many independent pieces, like images and the math behind neural networks.

In CUDA, the smallest worker is a thread. Threads are grouped into blocks, and all the blocks together are called a grid. When we start work on the GPU, we choose how many blocks we want and how many threads each block has.

Every thread runs the same code. So how does a thread know which piece of the work is its own? CUDA gives each thread a few built-in values. blockIdx.x is the number of its block, blockDim.x is how many threads a block has and threadIdx.x is its number inside the block. With a small formula, we turn them into one unique number for every thread. Hover over the threads below to see it. When it feels clear, press Quiz me and try to find the right thread yourself.

You may wonder why every name ends with .x. Blocks and threads can also be arranged in two or three dimensions, with .y and .z, which is handy for images and 3D data. Here we only use one dimension to keep things simple.

A function that we launch on the GPU like this is called a kernel. We mark it with __global__ and start it with the special <<<blocks, threads>>> syntax. Below is the classic first CUDA program: adding two arrays. Each thread adds only one pair of numbers, and the threads work in parallel.

There is a catch. Threads come in whole blocks, so we often start a few more threads than we need. If n is 10 and a block has 4 threads, we need 3 blocks, which means 12 threads. That is why each thread checks i < n before it writes. Real programs usually use bigger blocks, like 256 threads, but small numbers are easier to see. Change the number of elements and try turning off the bounds check.

__global__ void add(float *a, float *b, float *c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    c[i] = a[i] + b[i];
}
 
// launch
add<<<3, 4>>>(a, b, c, 10);

On most computers, the CPU and the GPU have separate memories. Before a kernel runs, the data is copied to the GPU, for example with cudaMemcpy, and the result is copied back at the end. These copies can take longer than the work itself. If we copy the data back and forth around every kernel, the GPU can spend most of its time waiting. If we keep the data on the GPU between kernels, we only pay for the copies once.

You meet this even if you never write CUDA yourself. In PyTorch, .to("cuda") copies a tensor into GPU memory, and .cpu() or .item() bring data back to the CPU. .item() also makes the CPU wait until the GPU has finished its work. That is why calling them again and again inside a training loop can slow it down: the data keeps moving back and forth, and the CPU keeps waiting.

Now we know all the pieces, so let's put them together. A small CUDA program has five steps: reserve memory on the GPU, copy the inputs, run the kernel, copy the result back and free the memory. Step through it below and watch both memories.

const int n = 10;
size_t size = n * sizeof(float);
float a[n], b[n], c[n]; // on the CPU, a and b are filled
float *d_a, *d_b, *d_c;          // will point to GPU memory
 
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
int threads = 256;
int blocks = (n + threads - 1) / threads;
add<<<blocks, threads>>>(d_a, d_b, d_c, n);
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

That is the main idea of CUDA. Split the work into many small pieces, let every thread find its own piece with a simple formula, and move the data as little as possible. Many GPU programs, from image filters to neural networks, are built on these same ideas.

@aykutkardas