We need to learn CUDA to increase the performance of models across domains world models (games, robotics, proteins), simulations, LLMs, and more. As models grow larger, they demand more compute. Currently, we are not able to fully utilize GPU power, which means energy and compute are being wasted. GPUs are the backbone of modern deep learning, so we must learn to optimize them as much as possible for our needs.

CUDA is all about parallel programming. The main goal of a GPU is to divide work among thousands of threads. When we write a GPU kernel, the primary goal is to maximize the throughput of a massive number of threads. There are a few challenges with GPU programming: if we use too few threads, the GPU performs poorly; designing parallel algorithms is in practice harder than designing sequential ones; and speed is often limited by memory latency and throughput, making many workloads memory bound rather than compute-bound.

Data parallelism refers to the phenomenon in which the computation work to be performed on different parts of the dataset can be done independently of each other and thus can be done in parallel with each other.

Images and videos have millions to billions of pixels. Scientific applications like fluid dynamics simulations work with billions of grid points. Molecular dynamics applications must simulate interactions between thousands to billions of atoms. Most of these pixels, particles, and grid points can usually be processed largely independently. For example, in image processing, converting a color pixel to grayscale requires only the data of that single pixel. Blurring an image by averaging each pixel's color with its neighbors can also be broken down into many smaller computations that execute independently. This kind of independent evaluation of different pieces of data is the foundation of data parallelism.

The structure of a CUDA C program reflects the coexistence of a CPU (host) and one or more GPUs (devices) in the same system. Each CUDA C source file can contain a mixture of host code and device code. By default, any traditional C program is a valid CUDA program that contains only host code. Device code can be added to any source file and is clearly marked with special CUDA C keywords. This device code includes functions called kernels, whose code is executed in a data parallel manner across many threads simultaneously.

Imagine you own a restaurant you (the head chef / CPU) read the recipe and decide what needs to be cooked. But you can’t chop 1,000 vegetables yourself so you shout the order to the kitchen floor, where hundreds of line cooks (GPU threads) are waiting. Every cook grabs one vegetable and chops it at the exact same moment. Done in seconds instead of hours.

That’s CUDA.

The execution rhythm :

CUDA kernel launches are asynchronous, meaning the CPU does not wait for the GPU to finish before continuing. The kernel is queued for execution and the CPU immediately moves on to the next instruction. Because of this, explicit synchronization is sometimes necessary specifically when the program depends on completed GPU results. Functions like cudaDeviceSynchronize() force the CPU to wait until all previously launched GPU work has finished, which is especially important before reading results back to the CPU or measuring execution time accurately.

CUDA functions do not fail silently they return error codes that should always be checked in real applications. Many issues like invalid memory access, illegal kernel launches, or misconfigured grid sizes only become visible through error checking. A common practice is to wrap CUDA calls with error-check macros or check cudaGetLastError() after kernel launches. This ensures that problems are detected early instead of producing incorrect results silently.

CUDA provides two levels of programming interfaces: the runtime API and the driver API. The runtime API (used in most examples like cudaMalloc, cudaMemcpy, and kernel launches) is higher-level and easier to use because it automatically manages context and setup. The driver API is lower-level, offering more control but requiring explicit management of GPU contexts and modules. Most application-level CUDA programming, especially in ML and simulation, is built using the runtime API.

GPU memory is organized in layers with different speeds and scopes. Registers are the fastest and are private to each thread. Shared memory is fast and shared among all threads within a block. Global memory is large but high-latency, and is accessible to all threads across all blocks. Efficient CUDA programs minimize global memory access and reuse data through shared memory or registers whenever possible. Understanding this hierarchy is critical because memory access patterns often determine performance more than raw compute throughput

The number of threads per block is usually chosen based on hardware efficiency rather than problem size alone. Common choices like 128, 256, or 512 threads align well with GPU warp size (32 threads per warp), allowing efficient scheduling. The goal is to keep enough active warps to hide memory latency while avoiding excessive resource usage per block, such as registers or shared memory. There is no single optimal value—it depends on the kernel and hardware.

__global__ the kernel qualifier

This marks a function as a kernel. The CPU launches it thousands of GPU threads run it in parallel.

// Only the CPU can call this. Thousands of GPU threads will run it.__global__ void myKernel(float* data, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) data[i] *= 2.0f;}

__device__ GPU helper functions

A function that runs on the GPU and is called by other GPU functions (kernels or other device functions). Useful for math utilities you want to reuse across kernels.

__device__ float square(float x) { return x * x;}

__global__ void squareKernel(float* data, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) data[i] = square(data[i]); // calling a __device__ function from a kernel}

CUDA C files use the .cu extension. You compile them with nvcc, NVIDIA’s compiler driver. It splits your code CPU portions go to your normal C++ compiler, GPU portions go to the PTX assembler.

nvcc --version # check CUDA toolkit versionnvidia-smi # check GPU and driver

Basic compile and run

nvcc vector_add.cu -o vector_add./vector_add

GPUs work best when data stays in GPU memory (VRAM) and many operations are queued back to back, rather than constantly moving data between the CPU and GPU.

Concrete example: In a AI Training loop, instead of doing:

forward → pull loss to CPU → check it → push back → backward

You want:

forward → backward → optimizer step → [all on GPU, no transfers]

Only pull to CPU occasionally (e.g., logging every N steps).

The cost of breaking this: Calling something like .item() on a PyTorch tensor mid computation forces a CPU-GPU sync, which can be a massive bottleneck this is a classic mistake in tight training loops.

In short: keep data on GPU, batch your kernel launches, and minimize round trips to CPU. okay back to our CUDA

Now let’s start with vector addition and write first cuda code. but before that let’s write one host code function works

void vecAdd(float* a_h, float* b_h, float* c_h, int n){ // Loop over all elements in the vectors for (int i = 0; i < n; i++){ c_h[i] = a_h[i] + b_h[i]; }}

int main(){ // Assume A, B, C are already allocated and initialized on CPU float* A; float* B; float* C; int N; // --------------------------------------------- // Preprocessing stage (CPU side) // We prepare / organize data (A, B, C, N) // before passing it to GPU kernel // --------------------------------------------- vecAdd(A, B, C, N); return 0;}

C program that consists of a main function and a vector addition function. In all our examples, whenever there is a need to distinguish between host and device data, we will suffix the names of variables that are used by the host with “_h”and those of variables that are used by a device with _d to remind ourselves of the intended usage of these variables. Since we have only host code in code_block, we see only variables suffixed with _h.

In current CUDA systems, devices are often hardware cards that come with their own dynamic random access memory called device global memory, or simply global memory. For example, an NVIDIA GTX 1650 which i’m currenly using comes with around 4GB of global memory. Calling it global memory distinguishes it from other types of device memory that are also accessible to programmers.

Your normal C program runs on the CPU, but CUDA kernels run on the GPU.Since the GPU has its own memory, data must be copied between them manually.

Think of it like this:

CPU (Host) RAM GPU (Device) VRAM----------------- -------------------A_h ----copy-----> A_dB_h ----copy-----> B_d

GPU kernel runs here

C_h <---copy----- C_d

cudaMalloc((void **) &A_d, size);cudaMalloc((void **) &B_d, size);cudaMalloc((void **) &C_d, size);

Reserve space inside GPU memory.

cudaMemcpy() --> memory data transfer it requires four parameters pointers to destination pointer to source, number of bytes copied, type/direction of transfer,cudaMemcpy(A_d, A_h, size, cudaMemcpyHostToDevice);cudaMemcpy(B_d, B_h, size, cudaMemcpyHostToDevice);cudaMemcpy(C_h, C_d, size, cudaMemcpyDeviceToHost);

Move inputs into GPU memory.

kernel<<<...>>>();

GPU processes data.

cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);

Bring results back.

cudaFree(...)

Release VRAM.

In order to allow memory for device we need to use cudaMalloc and we need to pass two parameters:

cudaFree() : Free object from device global memory , Pointer to freed object here is examples :

float *A_dint size = n * sizeof(float)cudaMalloc((void**)&A_d, size)cudaFree(A_d);

important things to remember is data transfer is expensive so we need to minimize copies keep data on gpu and if possible batch operations together

When a CUDA kernel launches thousands of threads, every thread executes the same kernel code. Because of this, each thread must be able to identify which block it belongs to, its position within that block, and which data element it is responsible for processing. CUDA provides built in variables for exactly this purpose: threadIdx, blockIdx, and blockDim.

These variables allow every thread to compute a unique global index.

The variable threadIdx gives each thread a unique coordinate inside its thread block.

For 1 dimensional thread organization, we use: threadIdx.x If a block contains 256 threads, the thread indices are: 0→255

So:

and so on these indices are local to the block. That means every block starts counting again from 0.

While threadIdx identifies a thread inside a block, blockIdx identifies the block itself.

For one-dimensional grids:

blockIdx.x

All threads inside the same block share the same blockIdx.x value.

Example:

CUDA also provides the block size through:

If each block contains 256 threads:

blockDim.x = 256

This value is identical for all threads in the same kernel launch.

Each thread combines:

to compute a unique global index.

The formula is:

i= blockIdx.x × blockDim.x + threadIdx.x

This is one of the most important formulas in CUDA programming.

Example with 256-thread blocks:

Threads compute:

Threads compute:

So every thread gets a unique position in the entire grid.

Each thread processes one vector element.

__global__void vecAddKernel(float* A, float* B, float* C, int n) {

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

if (i < n) { C[i] = A[i] + B[i]; }}

The kernel does not contain a traditional loop. instead, the grid of threads replaces the loop.

In CPU code:

for (int i = 0; i < n; i++) { C[i] = A[i] + B[i];}

In CUDA:

This is called loop parallelism.

Vector lengths are not always multiples of the block size.

Suppose:

We need 4 blocks:

4×32=128

But only 100 threads are needed.

That means:

The condition:

if (i < n)

prevents extra threads from accessing invalid memory.

The host launches the kernel using:

vecAddKernel<<<ceil(n / 256.0), 256>>>(A_d, B_d, C_d, n);

This means:

The grid automatically scales to large vectors.

Memory Flow

A complete CUDA program usually follows this sequence:

Example:

__global__ void vecAddKernel(float* A_d, float* B_d, float* C_d, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { C_d[i] = A_d[i] + B_d[i]; }}// ─── Host wrapper ─────// Handles all memory management and kernel launch.// A_h, B_h are inputs on CPU. Result is written into C_h.void vecAdd(float* A_h, float* B_h, float* C_h, int n) { float* A_d; float* B_d; float* C_d; int size = n * sizeof(float);

// Step 1: allocate GPU memory CUDA_CHECK(cudaMalloc((void**)&A_d, size)); CUDA_CHECK(cudaMalloc((void**)&B_d, size)); CUDA_CHECK(cudaMalloc((void**)&C_d, size));

// Step 2: copy inputs CPU → GPU CUDA_CHECK(cudaMemcpy(A_d, A_h, size, cudaMemcpyHostToDevice)); CUDA_CHECK(cudaMemcpy(B_d, B_h, size, cudaMemcpyHostToDevice));

// Step 3: launch kernel int threadsPerBlock = 256; int blocksPerGrid = (int)ceil((float)n / threadsPerBlock); vecAddKernel<<<blocksPerGrid, threadsPerBlock>>>(A_d, B_d, C_d, n); CUDA_CHECK(cudaGetLastError()); // catch bad launch config CUDA_CHECK(cudaDeviceSynchronize()); // wait for kernel to finish

// Step 4: copy result GPU → CPU CUDA_CHECK(cudaMemcpy(C_h, C_d, size, cudaMemcpyDeviceToHost));

// Step 5: free GPU memory CUDA_CHECK(cudaFree(A_d)); CUDA_CHECK(cudaFree(B_d)); CUDA_CHECK(cudaFree(C_d));}

// ─── Main ────int main() { int n = 1000; float A_h[1000], B_h[1000], C_h[1000];

// Initialize: A[i] = i, B[i] = n - i // Expected result: C[i] = n = 1000 for all i for (int i = 0; i < n; i++) { A_h[i] = (float)i; B_h[i] = (float)(n - i); }

vecAdd(A_h, B_h, C_h, n);

// Verify a few results printf("C[0] = %.1f (expected 1000.0)\n", C_h[0]); printf("C[500] = %.1f (expected 1000.0)\n", C_h[500]); printf("C[999] = %.1f (expected 1000.0)\n", C_h[999]);

return 0;}