Cuda kernel basics

CUDA kernel basics

A kernel launches one function across thousands of threads, each working on its own slice of data. The classic first attempt is correct but an order of magnitude too slow: it indexes memory like a serial loop and sizes thread blocks by guesswork. These habits, applied on every kernel, close most of that gap before a profiler is even involved.

Method

1. Decompose the work onto grid, block, and thread. Compute the global index as int i = blockIdx.x * blockDim.x + threadIdx.x and bound it with if (i >= n) return; so the last, partly full block does not read past the array. Keep this one-thread-per-element version as the reference you optimize against.
2. Write a grid-stride loop for sizes larger than the grid. Loop for (int i = tid; i < n; i += blockDim.x * gridDim.x) so a fixed grid handles any n and each thread covers several elements. This decouples the launch configuration from problem size and keeps blocks busy.
3. Size blocks in multiples of the 32-thread warp. Threads issue in warps of 32, so a block of 200 wastes 24 lanes in its final warp. Start at 128 or 256 threads and compute the grid as (n + block - 1) / block. Never launch a block dimension that is not a multiple of 32.
4. Coalesce global loads and stores. When the 32 threads of a warp touch consecutive addresses, the hardware fuses them into one 128-byte transaction. a[i] on the global index coalesces; a[i * stride] or indexing the outer axis of a 2D array scatters into up to 32 separate sectors. Make threadIdx.x walk the innermost, contiguous axis.
5. Check occupancy with the tools, not by feel. Call cudaOccupancyMaxPotentialBlockSize or read the occupancy line in Nsight Compute. Registers per thread and shared memory per block cap how many warps a streaming multiprocessor keeps resident; overspending either starves the SM of work to hide memory latency. Target at least 50 percent theoretical occupancy before blaming the algorithm.
6. Prove correctness before timing anything. Wrap launches with cudaGetLastError() and a cudaDeviceSynchronize() in debug builds, and run once under compute-sanitizer to catch out-of-bounds writes and races that corrupt results only on some launches.

Checks

• Does compute-sanitizer ./app report zero errors on a representative input?
• Is every block dimension a multiple of 32, and does the grid cover the tail?
• Does Nsight Compute report global accesses as coalesced (high sectors-per-request efficiency) rather than scattered?
• Does the kernel match a trusted CPU reference within floating-point tolerance?

Boundaries

This is single-kernel correctness and first-order speed on one GPU. It does not cover streams and copy overlap, multi-GPU NCCL collectives, or wringing the last percent from a GEMM. Once loads coalesce and occupancy is healthy, move to a profiler-driven pass (kernel-profiling-nsight) and to staging data across the memory tiers (gpu-memory-hierarchy) instead of guessing.