CUDA C++ Kernel Anatomy

A CUDA C++ program has two halves: host code that runs on the CPU, allocates GPU memory and launches work, and device code (kernels) that runs on the GPU, one copy per thread. Every program below ran on a Tesla T4 with CUDA 13.0, built with plain nvcc.

One thread per element

A kernel is a __global__ function. The launch <<<blocks, threads>>> decides how many copies run, and each copy works out which element is its own.

// vecadd.cu - copy in, launch one thread per element, copy out
#include <cstdio>
#include <cuda_runtime.h>

__global__ void vec_add(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;  // global index of this thread
  if (i < n) c[i] = a[i] + b[i];                  // spare threads in the last block do nothing
}

int main() {
  const int n = 1000000;
  const size_t bytes = n * sizeof(float);
  float *ha = new float[n], *hb = new float[n], *hc = new float[n];
  for (int i = 0; i < n; i++) { ha[i] = i; hb[i] = 2.0f * i; }
  float *da, *db, *dc;
  cudaMalloc(&da, bytes); cudaMalloc(&db, bytes); cudaMalloc(&dc, bytes);
  cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice);
  cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice);
  int threads = 256, blocks = (n + threads - 1) / threads;
  vec_add<<<blocks, threads>>>(da, db, dc, n);
  cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost);  // waits for the kernel to finish
  printf("%d blocks x %d threads = %d threads for %d elements\n", blocks, threads, blocks * threads, n);
  printf("c[1] = %.0f, c[999999] = %.0f\n", hc[1], hc[n - 1]);
  cudaFree(da); cudaFree(db); cudaFree(dc);
}

$ nvcc -o vecadd vecadd.cu && ./vecadd
3907 blocks x 256 threads = 1000192 threads for 1000000 elements
c[1] = 3, c[999999] = 2999997

blockIdx.x * blockDim.x + threadIdx.x turns “thread 3 of block 2” into one global number, so each of the million threads adds one pair. The block count is rounded up so every element gets a thread: 3907 blocks of 256 is 1,000,192 threads, 192 more than there are elements. Those spare threads land in the last block, and the if (i < n) guard is what stops them writing past the end of c.

The kernel can’t touch ha or hb, because they live in CPU memory. That’s why the data is copied to da and db first and the result copied back. The launch itself returns at once; it’s the final cudaMemcpy that waits for the kernel to finish before copying. c[999999] is 999,999 + 1,999,998, so the last element was computed too.

A launch that fails without a word

A launch doesn’t return an error code. Ask for more threads per block than the GPU allows, and nothing tells you unless you check.

// errors.cu - a bad launch fails silently unless you ask
#include <cstdio>
#include <cuda_runtime.h>

__global__ void fill(float* x, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) x[i] = 1.0f;
}

int main() {
  const int n = 4096;
  float* x; cudaMalloc(&x, n * sizeof(float));
  fill<<<n / 2048, 2048>>>(x, n);            // 2048 threads per block is over the limit
  printf("the bad launch returned, and nothing was printed\n");
  cudaError_t err = cudaGetLastError();      // launch errors are reported here
  printf("cudaGetLastError: %s\n", cudaGetErrorString(err));
  fill<<<n / 256, 256>>>(x, n);
  printf("with 256 per block: %s\n", cudaGetErrorString(cudaGetLastError()));
  printf("cudaDeviceSynchronize: %s\n", cudaGetErrorString(cudaDeviceSynchronize()));
  cudaFree(x);
}

$ nvcc -o errors errors.cu && ./errors
the bad launch returned, and nothing was printed
cudaGetLastError: invalid argument
with 256 per block: no error
cudaDeviceSynchronize: no error

The first launch asked for 2048 threads per block, and the T4 allows 1024. The kernel never ran, yet the program carried on and printed its next line as if nothing had happened. The failure only surfaced when cudaGetLastError() was called; CUDA 13.0 words it as invalid argument.

cudaGetLastError() also resets the error, which is why the check after the second, valid launch says no error. cudaDeviceSynchronize() waits for all launched work and returns any error raised while kernels ran; here there was none, because the bad kernel never started. In real code, wrap every CUDA call and check after every launch; Practice 2 builds a macro for it.

Where each function may run

__global__ is one of three qualifiers. They tell nvcc which processor each function is compiled for.

// qualifiers.cu - where each function is allowed to run
#include <cstdio>
#include <cuda_runtime.h>

__host__ __device__ float square(float v) { return v * v; }  // compiled for CPU and GPU
__device__ float plus_one(float v) { return v + 1.0f; }      // GPU only, called from kernels

__global__ void kernel(float* out) {
  int i = threadIdx.x;
  out[i] = plus_one(square(i));
  if (i == 3) printf("device: thread %d computed %.0f\n", i, out[i]);
}

int main() {
  float* out; cudaMallocManaged(&out, 4 * sizeof(float));  // one pointer for CPU and GPU
  kernel<<<1, 4>>>(out);
  printf("host: launched, the kernel may still be running\n");
  cudaDeviceSynchronize();
  printf("host: square(3) on the CPU = %.0f\n", square(3.0f));
  printf("host: out = %.0f %.0f %.0f %.0f\n", out[0], out[1], out[2], out[3]);
  cudaFree(out);
}

$ nvcc -o qualifiers qualifiers.cu && ./qualifiers
host: launched, the kernel may still be running
device: thread 3 computed 10
host: square(3) on the CPU = 9
host: out = 1 2 5 10

square is __host__ __device__, so it’s compiled twice: the kernel calls the GPU copy, and main calls the CPU copy. plus_one is __device__ only, so kernels can call it and the CPU can’t. Each thread computed i * i + 1, giving 1, 2, 5 and 10.

The order of the lines shows the launch is asynchronous. The host’s “launched” line printed first, even though the kernel was launched before it. printf inside a kernel writes to a buffer that is only flushed to the screen when the host synchronizes, so the device line appeared at cudaDeviceSynchronize(). cudaMallocManaged gives one pointer that both sides can use, which is why main can read out directly after the sync, with no copy.

A 2D grid for 2D data

Images and matrices map more naturally onto a 2D launch. dim3 gives a block and a grid an x and a y size.

// grid2d.cu - a 2D grid of 2D blocks over a 1000 x 600 image
#include <cstdio>
#include <cuda_runtime.h>

__global__ void brighten(unsigned char* img, int w, int h) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;  // column
  int y = blockIdx.y * blockDim.y + threadIdx.y;  // row
  if (x < w && y < h) img[y * w + x] = min(255, img[y * w + x] + 50);
}

int main() {
  const int w = 1000, h = 600;
  unsigned char* img; cudaMallocManaged(&img, w * h);
  for (int i = 0; i < w * h; i++) img[i] = i % 256;
  dim3 block(16, 16);
  dim3 grid((w + block.x - 1) / block.x, (h + block.y - 1) / block.y);
  brighten<<<grid, block>>>(img, w, h);
  cudaDeviceSynchronize();
  printf("grid %u x %u blocks of %u x %u threads = %u threads for %d pixels\n",
         grid.x, grid.y, block.x, block.y, grid.x * grid.y * block.x * block.y, w * h);
  int x = 999, y = 599;
  printf("pixel (%d, %d) went from %d to %d\n", x, y, (y * w + x) % 256, img[y * w + x]);
  cudaFree(img);
}

$ nvcc -o grid2d grid2d.cu && ./grid2d
grid 63 x 38 blocks of 16 x 16 threads = 612864 threads for 600000 pixels
pixel (999, 599) went from 191 to 241

Each thread gets a column from the x parts and a row from the y parts, then flattens them into y * w + x, the row-major index of its pixel. The grid is rounded up in both directions: 1000 / 16 is 62.5, so 63 blocks across, and 600 / 16 is 37.5, so 38 blocks down. That’s 612,864 threads for 600,000 pixels, and the guard now has to check both x < w and y < h.

The bottom-right pixel started at 599,999 mod 256 = 191 and came out 50 brighter, at 241. min caps the value at 255; without it, a bright pixel would wrap round past 255 to a dark one.

Common gotchas

• No bounds guard. The grid is rounded up to whole blocks, so the last block almost always has threads with no element. Without if (i < n) they write past the array.
• Not checking errors. A launch with a bad configuration simply doesn’t run. Call cudaGetLastError() after launches and check the return value of every CUDA call.
• Reading results before the kernel finishes. Launches return at once. A cudaMemcpy back to the host waits for you; with managed memory you must call cudaDeviceSynchronize() yourself.
• Mixing up x and y. In y * w + x, x is the column and y the row. Swap them and a non-square image is processed in the wrong order, or out of bounds.

Practice

1. Print the launch limits of your GPU: its name and compute capability, the maximum threads per block, the maximum block size in each dimension, the maximum grid width, the warp size and the number of SMs.

// limits.cu - the launch limits this GPU enforces
#include <cstdio>
#include <cuda_runtime.h>

int main() {
  cudaDeviceProp p;
  cudaGetDeviceProperties(&p, 0);
  printf("%s, compute capability %d.%d\n", p.name, p.major, p.minor);
  printf("max threads per block: %d\n", p.maxThreadsPerBlock);
  printf("max block size (x, y, z): %d, %d, %d\n", p.maxThreadsDim[0], p.maxThreadsDim[1], p.maxThreadsDim[2]);
  printf("max grid size x: %d\n", p.maxGridSize[0]);
  printf("warp size: %d, SMs: %d\n", p.warpSize, p.multiProcessorCount);
}

$ nvcc -o limits limits.cu && ./limits
Tesla T4, compute capability 7.5
max threads per block: 1024
max block size (x, y, z): 1024, 1024, 64
max grid size x: 2147483647
warp size: 32, SMs: 40

2. Write a CHECK macro that prints the file, the line and the error string of a failed CUDA call and stops the program. Use it on the cudaMalloc and after both launches from errors.cu.

// check.cu - a CHECK macro that names the file and line of a failed call
#include <cstdio>
#include <cuda_runtime.h>

#define CHECK(call) do { cudaError_t e = (call); if (e != cudaSuccess) { \
  printf("%s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(e)); return 1; } } while (0)

__global__ void fill(float* x, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) x[i] = 1.0f;
}

int main() {
  const int n = 4096;
  float* x;
  CHECK(cudaMalloc(&x, n * sizeof(float)));
  fill<<<n / 256, 256>>>(x, n);
  CHECK(cudaGetLastError());
  printf("256 per block: ok\n");
  fill<<<n / 2048, 2048>>>(x, n);
  CHECK(cudaGetLastError());
  printf("never printed\n");
}

$ nvcc -o check check.cu && ./check
256 per block: ok
check.cu:21: invalid argument

3. Call the __device__ function plus_one directly from main. What does nvcc say?

// hostcall.cu - calling a __device__ function from the CPU
#include <cstdio>

__device__ float plus_one(float v) { return v + 1.0f; }

int main() {
  printf("%.0f\n", plus_one(2.0f));
}

$ nvcc -o hostcall hostcall.cu
hostcall.cu(7): error: calling a __device__ function("plus_one(float)") from a __host__ function("main") is not allowed
    printf("%.0f\n", plus_one(2.0f));
                     ^

1 error detected in the compilation of "hostcall.cu".

4. Change the image kernel to use 32 x 8 blocks. How big is the grid now, and does every pixel still get processed? Count the pixels that changed to find out.

// grid32x8.cu - the same image with wide 32 x 8 blocks
#include <cstdio>
#include <cuda_runtime.h>

__global__ void brighten(unsigned char* img, int w, int h) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;
  if (x < w && y < h) img[y * w + x] = min(255, img[y * w + x] + 50);
}

int main() {
  const int w = 1000, h = 600;
  unsigned char* img; cudaMallocManaged(&img, w * h);
  for (int i = 0; i < w * h; i++) img[i] = i % 256;
  dim3 block(32, 8);
  dim3 grid((w + block.x - 1) / block.x, (h + block.y - 1) / block.y);
  brighten<<<grid, block>>>(img, w, h);
  cudaDeviceSynchronize();
  printf("grid %u x %u blocks of %u x %u threads = %u threads for %d pixels\n",
         grid.x, grid.y, block.x, block.y, grid.x * grid.y * block.x * block.y, w * h);
  int changed = 0;
  for (int i = 0; i < w * h; i++) changed += img[i] != i % 256;
  printf("pixels changed: %d\n", changed);
  cudaFree(img);
}

$ nvcc -o grid32x8 grid32x8.cu && ./grid32x8
grid 32 x 75 blocks of 32 x 8 threads = 614400 threads for 600000 pixels
pixels changed: 597657