CUDA

Learning Objectives

  • Explain the host/device memory model and why explicit data transfers exist

  • Write a CUDA kernel and configure its grid and block dimensions

  • Allocate GPU memory with cudaMalloc and transfer data with cudaMemcpy

  • Check CUDA API error codes and kernel launch errors

  • Measure kernel execution time separately from data transfer time

  • Compute effective memory bandwidth and compare to GPU specs

Note

If you’re just beginning, it’s highly recommended that you use Kokkos instead. It’s programmed largely the same as cuda but acts as a portability layer that lets your code run on GPUs other than nvidia.

GPU Execution Model

A GPU consists of thousands of simple cores organized into Streaming Multiprocessors (SMs). Unlike a CPU, which has a small number of powerful, latency-optimized cores, a GPU is designed to hide memory latency by switching between thousands of in-flight threads.

When you launch a CUDA kernel, you specify:

  • A grid — the number of thread blocks to launch

  • A thread block — the number of threads within each block

Every thread executes the same kernel function but with a unique index (SIMT: Single Instruction, Multiple Thread).

Thread Indexing (1D case)

global_thread_index = blockIdx.x * blockDim.x + threadIdx.x

Variable

Meaning

threadIdx.x

Thread index within its block (0 to blockDim.x - 1)

blockIdx.x

Block index within the grid

blockDim.x

Number of threads per block

gridDim.x

Number of blocks in the grid

Block size must be a multiple of 32 (the warp size). Common choices: 256 or 512 threads per block.

Warning

When n is not evenly divisible by blocksize, the last block contains threads whose computed index exceeds the array bounds. Always guard with if (i >= n) return; at the start of the kernel.

Host and Device Memory

CPUs and GPUs have separate memory spaces. The GPU has its own DRAM (device memory) connected via a high-bandwidth bus (900–3000 GB/s for HBM3). CPU-to-GPU transfers go over PCIe (~16 GB/s) — a significant bottleneck.

The standard workflow:

  1. Allocate host (CPU) arrays with malloc

  2. Allocate device (GPU) arrays with cudaMalloc

  3. Copy input data host → device with cudaMemcpy(..., cudaMemcpyHostToDevice)

  4. Launch kernel(s)

  5. cudaDeviceSynchronize() — wait for the GPU to finish

  6. Copy results device → host with cudaMemcpy(..., cudaMemcpyDeviceToHost)

  7. cudaFree device arrays; free host arrays

Tip

A common convention is to suffix device pointers with _d (e.g., a_d, b_d) to distinguish them from host pointers.

CUDA Stream Triad Kernel

The Stream Triad computes:

\[c_i = a_i + s \cdot b_i \quad \forall \, i \in [0, N)\]
__global__ void StreamTriad(const int n, const double scalar,
                            const double *a, const double *b, double *c) {
  const int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= n) return;
  c[i] = a[i] + scalar * b[i];
}

int main() {
  const int n = 80000000;
  double *a = (double *)malloc(n * sizeof(double));
  double *b = (double *)malloc(n * sizeof(double));
  double *c = (double *)malloc(n * sizeof(double));
  
  // Initialize host arrays
  for (int i = 0; i < n; i++) { a[i] = 1.0; b[i] = 2.0; }

  double *a_d, *b_d, *c_d;
  cudaMalloc(&a_d, n * sizeof(double));
  cudaMalloc(&b_d, n * sizeof(double));
  cudaMalloc(&c_d, n * sizeof(double));

  const int blocksize = 512;
  const int gridsize = (n + blocksize - 1) / blocksize;

  // Copy to device
  cudaMemcpy(a_d, a, n * sizeof(double), cudaMemcpyHostToDevice);
  cudaMemcpy(b_d, b, n * sizeof(double), cudaMemcpyHostToDevice);

  // Launch kernel
  StreamTriad<<<gridsize, blocksize>>>(n, 3.0, a_d, b_d, c_d);
  cudaDeviceSynchronize();

  // Copy back
  cudaMemcpy(c, c_d, n * sizeof(double), cudaMemcpyDeviceToHost);

  cudaFree(a_d); cudaFree(b_d); cudaFree(c_d);
  free(a); free(b); free(c);
}

Building and Running

On Workstation

nvcc -O3 stream_triad.cu -o stream_triad
./stream_triad

On ARC Clusters with GPU hardware

  1. Load CUDA module:

    module load cuda

  2. Request GPU in SLURM:

    #SBATCH --gres=gpu:1
#SBATCH --partition=<valid_gpu_partition>

  3. Submit:

    sbatch run.slurm

See GPU Usage for cluster-specific details.

Performance Measurement

The effective bandwidth for the Stream Triad is:

\[B = \frac{3 \times N \times 8 \text{ bytes}}{t \text{ seconds}} \text{ GB/s}\]

where \(N\) is array size and \(t\) is runtime. Compare to GPU specs:

GPU

Memory Bandwidth

A100 80GB

~2000 GB/s

V100 32GB

~900 GB/s

L40S

~864 GB/s

References