CUDA C++ Kernel Anatomy

What You Need to Know (Theory)

CUDA C++ adds NVIDIA extensions on top of standard C++:

• Host code runs on CPU (allocations, kernel launches).
• Device code (kernels) runs on GPU.

Kernel launch configuration

A launch like:

kernel<<<grid, block>>>(args...);

means:

• grid = number of blocks
• block = number of threads per block

Memory access qualifiers

Qualifiers help the compiler understand usage:

• const T* often maps to read-only behavior (and can enable cache paths)
• __restrict__ can help optimization by telling the compiler pointers don’t alias

Code Example 1 — Vector Add in CUDA C++

// vector_add.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void vector_add(const float* __restrict__ A,
                            const float* __restrict__ 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;
  const size_t bytes = N * sizeof(float);

  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;
  }

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

  cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

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

  cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

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

  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 Image-like Kernel

Many GPU tasks are 2D (images, matrices). Use threadIdx.{x,y} and blockIdx.{x,y}.

#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; // col
  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 = 512, H = 512;
  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

• Forgetting bounds checks (if (i < N) / if (x < W && y < H)).
• Mismatched indexing between host-side assumptions and kernel-side layout.
• Using the wrong qualifiers: don’t add const/__restrict__ unless correct—wrong assumptions can break correctness.

Quick Checklist

• Compute global indices correctly (blockIdx*blockDim + threadIdx)
• Choose a block size that matches your problem dimensions
• Guard against out-of-range threads
• Validate results on small inputs before scaling up