CUDA Programming Basics

What You’re Building

CUDA is NVIDIA’s platform for GPU computing. You write kernels (GPU functions), and launch them with a grid/block/thread configuration.

A classic first program is vector addition: C[i] = A[i] + B[i].

Core CUDA Concepts (Theory)

1) Thread hierarchy

• grid: the whole collection of thread blocks
• block: a group of threads that can cooperate using shared memory
• thread: the smallest execution unit running the kernel

2) Built-in indices

Inside a kernel, CUDA provides built-in variables:

• blockIdx.x, blockDim.x, threadIdx.x

Common pattern:

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

3) Device vs Host memory

• Host (CPU) memory: normal allocations (malloc, new)
• Device (GPU) memory: allocated with cudaMalloc
• Move data using cudaMemcpy

Code Example 1 — Vector Add (Single Kernel)

vector_add.cu

#include <cstdio>
#include <cuda_runtime.h>

__global__ void vector_add(const float* A, const float* B, float* C, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x; // global thread index
  if (i < N) {
    C[i] = A[i] + B[i];
  }
}

int main() {
  const int N = 1 << 20; // ~1 million
  const size_t bytes = N * sizeof(float);

  // 1) Host allocations
  float* h_A = (float*)malloc(bytes);
  float* h_B = (float*)malloc(bytes);
  float* h_C = (float*)malloc(bytes);

  for (int i = 0; i < N; i++) {
    h_A[i] = 1.0f;
    h_B[i] = 2.0f;
  }

  // 2) Device allocations
  float *d_A, *d_B, *d_C;
  cudaMalloc(&d_A, bytes);
  cudaMalloc(&d_B, bytes);
  cudaMalloc(&d_C, bytes);

  // 3) Copy host -> device
  cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

  // 4) Launch kernel
  int threads = 256;
  int blocks = (N + threads - 1) / threads;
  vector_add<<<blocks, threads>>>(d_A, d_B, d_C, N);

  // 5) Copy device -> host
  cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

  // 6) Verify
  printf("C[0]=%f  C[N-1]=%f\n", h_C[0], h_C[N-1]);

  // Cleanup
  cudaFree(d_A);
  cudaFree(d_B);
  cudaFree(d_C);
  free(h_A);
  free(h_B);
  free(h_C);
  return 0;
}

Compile & Run

nvcc -O2 -arch=sm_75 vector_add.cu -o vector_add
./vector_add

Code Example 2 — 2D Grid for Matrix-like Work

Many problems map naturally to 2D:

• rows ↔ threadIdx.y
• cols ↔ threadIdx.x

#include <cuda_runtime.h>
#include <cstdio>

__global__ void add_2d(const float* A, const float* B, float* C, 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) {
    int idx = y * W + x;
    C[idx] = A[idx] + B[idx];
  }
}

int main() {
  int W = 1024, H = 1024;
  size_t bytes = W * H * sizeof(float);

  float *d_A, *d_B, *d_C;
  cudaMalloc(&d_A, bytes);
  cudaMalloc(&d_B, bytes);
  cudaMalloc(&d_C, bytes);

  dim3 threads(16, 16);
  dim3 blocks((W + threads.x - 1) / threads.x,
              (H + threads.y - 1) / threads.y);

  add_2d<<<blocks, threads>>>(d_A, d_B, d_C, W, H);
  cudaFree(d_A);
  cudaFree(d_B);
  cudaFree(d_C);
  return 0;
}

Common Gotchas

• Out-of-bounds memory: always check if (i < N) (or x < W && y < H).
• Choosing block size:
  • too small → low occupancy
  • too large → register/shared memory pressure
• Kernel launch errors: check errors in real projects using cudaGetLastError() and cudaDeviceSynchronize() after launches.

Quick Checklist

• Allocate device buffers with cudaMalloc
• Copy inputs to device with cudaMemcpy(..., HostToDevice)
• Launch kernels with correct grid/block sizes
• Copy outputs back to host with cudaMemcpy(..., DeviceToHost)