CUDA Programming Basics: GPU Computing Fundamentals

One frame, sixteen million pixels, no time to wait

A ground station receiving a Cartosat-class panchromatic image after a flood needs to run edge detection on the frame before the next pass overwrites the buffer — change analysis over a disaster zone cannot wait for a slow pipeline. The image is 4096×4096 pixels, which is exactly 16,777,216 pixels (4096² = 2¹² squared = 2²&sup4;). Edge detection here means a 3×3 Sobel convolution: for every interior pixel, compute two weighted sums (G_x and G_y) over its eight neighbours, then combine them into a gradient magnitude. Crucially, the value written at pixel (120, 340) depends only on the 3×3 neighbourhood around (120, 340) — never on the result at any other pixel. Every one of the 16.7 million outputs is computable independently of every other one.

That single property — independent, identical work repeated over millions of data elements — is exactly what a GPU is built for and what a CPU is not. A CPU core advances one instruction stream (or a handful, with wide SIMD) at a time; to process 16.7 million independent pixels, it must loop through them, one iteration after another, however cleverly the compiler pipelines that loop. A GPU instead assigns one thread to one pixel and runs enormous numbers of those threads at once. This chapter builds the CUDA programming model from first principles — the abstractions that let a programmer say "run this function once per pixel" and have the hardware actually do it in parallel — using this Sobel pass as the running example.

Why a GPU is shaped the way it is

NVIDIA's engineers documented the reasoning formally: a general-purpose graphics processor trades a small number of powerful, latency-optimized cores (a CPU's design point) for a very large number of simpler, throughput-optimized cores, because graphics — and, it turned out, a wide class of numerical workloads — is dominated by exactly the kind of massively data-parallel, arithmetic-heavy work a Sobel filter represents (Lindholm, Nickolls, Oberman & Montrym, "NVIDIA Tesla: A Unified Graphics and Computing Architecture," IEEE Micro, vol. 28, no. 2, 2008, pp. 39–55). NVIDIA released CUDA in 2007 specifically to expose that hardware to general C/C++ programmers, not just to graphics APIs.

Concretely, on an NVIDIA A100 data-center GPU there are 108 streaming multiprocessors (SMs), each capable of holding up to 2,048 threads resident at once. That is 108 × 2,048 = 221,184 threads physically in flight simultaneously on one chip — not queued, not time-sliced one at a time, but actually occupying hardware execution slots together. A CPU core, even a wide superscalar one, is not built to hold six-figure counts of independent instruction streams in flight; a GPU is built for nothing else. This is the architectural fact the rest of this chapter turns into a programming model.

Host and device: two separate computers talking

CUDA's central idea is that a program has two halves running on two separate processors with two separate memories: the host (the CPU, with system RAM) and the device (the GPU, with its own DRAM, typically HBM2 or HBM2e on data-center cards). Ordinary C/C++ code runs on the host. A function marked __global__ is a kernel — it runs on the device, and the host launches it.

Because host and device memory are physically separate, data has to be explicitly copied across the PCIe or NVLink bus before a kernel can touch it, and copied back afterward to read the result. cudaMalloc allocates device memory; cudaMemcpy moves bytes between host and device. A kernel launch itself does not return a result — it schedules work on the device and (by default) returns control to the host immediately, so the subsequent cudaMemcpy back to the host also acts as a synchronization point, blocking until the kernel has finished writing.

Grid, block, thread: the shape of parallel work

A kernel launch does not create "however many threads happen to be needed" as a flat pool. CUDA organizes threads into a fixed three-level hierarchy:

• Thread — the smallest unit of execution; runs one instance of the kernel body, with private registers.
• Block — a group of threads (up to 1,024 on current hardware) that run on the same SM, can synchronize with each other via __syncthreads(), and share a fast on-chip memory region.
• Grid — all the blocks launched by one kernel call, covering the entire problem.

The launch syntax kernel<<<blocksPerGrid, threadsPerBlock>>>(args) fixes the grid's and each block's dimensions. Inside the kernel, four built-in variables tell each thread who it is: threadIdx (position within its block), blockIdx (position of its block within the grid), blockDim (size of a block), and gridDim (size of the grid). For the common one-dimensional case, every thread computes its own unique position in the overall problem with one line:

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

This is the single most important line in CUDA programming — it is how "run once per data element" turns into "each thread knows which element is mine." The diagram below traces this formula for the fully worked example that follows.

Worked example: vector addition, fully traced

Before tackling a 2D image, work the exact indexing arithmetic for the one-dimensional case, because every 2D and 3D kernel is built from the same reasoning. Take two arrays of N = 1,000,000 floats and add them element-wise on the GPU.

#include <cuda_runtime.h>
#include <stdlib.h>

__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int N = 1000000;
    size_t size = N * sizeof(float);

    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    // fill h_A and h_B with input values here

    float *d_A, *d_B, *d_C;
    cudaMalloc((void**)&d_A, size);
    cudaMalloc((void**)&d_B, size);
    cudaMalloc((void**)&d_C, size);

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

    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

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

Trace the launch configuration by hand. blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock = (1,000,000 + 255) / 256 = 1,000,255 / 256. Integer division: 256 × 3906 = 999,936, and 256 × 3907 = 1,000,192; since 999,936 ≤ 1,000,255 < 1,000,448, the quotient is 3,907 blocks — this is the standard "ceiling division" trick for turning an element count that doesn't divide evenly into a whole number of blocks.

Total threads actually launched: 3,907 × 256 = 1,000,192 — 192 more than N. Trace the last block, blockIdx.x = 3906: its offset is 3906 × 256 = 999,936. For threadIdx.x = 63, idx = 999,936 + 63 = 999,999, which is exactly N - 1, the last valid element — still inside the array. For threadIdx.x = 64, idx = 999,936 + 64 = 1,000,000 = N, which is one past the end. Without the guard if (idx < N), that thread would write past the allocated buffer. In this last block, threads 0–63 (64 threads) are valid and threads 64–255 (192 threads) are masked off — matching the 192 "extra" threads counted above exactly. This is why the boundary check is not optional defensive style; without it, this exact kernel corrupts adjacent device memory on every launch where N is not a multiple of threadsPerBlock.

From 1D to 2D: the same formula, one axis each

The Sobel kernel from the opening scenario needs a row and a column, so the launch configuration and the index formula simply get a second axis, using CUDA's dim3 type instead of a plain integer:

__global__ void sobelEdge(const unsigned char *in, unsigned char *out,
                           int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= 1 && x < width - 1 && y >= 1 && y < height - 1) {
        int gx = -in[(y-1)*width+(x-1)] + in[(y-1)*width+(x+1)]
                 - 2*in[y*width+(x-1)]  + 2*in[y*width+(x+1)]
                 - in[(y+1)*width+(x-1)] + in[(y+1)*width+(x+1)];
        int gy = -in[(y-1)*width+(x-1)] - 2*in[(y-1)*width+x] - in[(y-1)*width+(x+1)]
                 + in[(y+1)*width+(x-1)] + 2*in[(y+1)*width+x] + in[(y+1)*width+(x+1)];
        int mag = abs(gx) + abs(gy);
        out[y*width+x] = (mag > 255) ? 255 : mag;
    }
}

// host launch configuration (mirrors the cudaMalloc/cudaMemcpy pattern above)
dim3 threadsPerBlock2D(16, 16);
dim3 blocksPerGrid2D((width + 15) / 16, (height + 15) / 16);
sobelEdge<<<blocksPerGrid2D, threadsPerBlock2D>>>(d_in, d_out, width, height);

With width = height = 4096, blocksPerGrid2D = (256, 256), giving 256 × 16 × 256 × 16 = 16,777,216 threads — one per pixel, exactly, with no boundary waste this time because 4096 divides evenly by 16. Each thread does about 17 arithmetic operations: computing gx combines six signed terms (five add/subtract steps) plus two multiplications by 2, the same for gy, plus two abs() calls and one add for the magnitude — 7 + 7 + 3 = 17. Across the whole frame that is 16,777,216 × 17 = 285,212,672, roughly 285 million arithmetic operations, all independent, all expressible as "one thread, one pixel." Spread across an A100's 221,184 simultaneously resident thread slots, covering all 16,777,216 pixels takes about 16,777,216 / 221,184 ≈ 76 waves of full-occupancy execution — versus a single CPU thread working through the same 16.7 million pixels one iteration at a time.

Memory: not one pool, three, with very different costs

The diagram already separates two of the three memory kinds a CUDA program deals with. Global memory is the device's main DRAM (HBM2 on an A100), visible to every thread in every block, with high aggregate bandwidth — 1,555 GB/s on the A100 40GB card, 2,039 GB/s on the 80GB card — but every access still costs on the order of hundreds of clock cycles of latency, because it is physically off the SM chip. Shared memory is on-chip SRAM, private to one block (an A100 SM offers up to 192 KB combined for shared memory and L1 cache), and is roughly one to two orders of magnitude lower latency than global memory. It exists specifically so threads within a block can cooperate: for example, a tiled matrix-multiply kernel loads a tile of the input matrices into shared memory once, then has many threads reuse it, instead of every thread re-fetching the same values from slow global memory. Registers are the fastest tier of all — private to a single thread, holding its loop counters and intermediate values like gx, gy, and mag in the Sobel kernel — but they are a scarce, per-SM resource; a kernel that uses too many registers per thread reduces how many threads can be resident at once.

The misconception: "each thread runs independently, like a CPU thread"

The single most common wrong mental model students bring to CUDA is treating a thread the way they think of an operating-system thread: an independent stream of execution that the hardware schedules whenever it likes, with branches costing nothing beyond the branch itself. That model is wrong at the level CUDA actually executes at. Threads are scheduled in groups of 32 called a warp, and within a warp, execution is SIMT (single instruction, multiple threads): all 32 lanes of a warp execute the same instruction at the same clock, in lockstep, with a per-lane active/inactive mask. If threads in the same warp take different branches of an if/else, the hardware does not run them concurrently — it runs the if path with the else-taking lanes masked idle, then runs the else path with the if-taking lanes masked idle. This is called warp divergence, and it means the warp's total time is the sum of both paths' costs, not the maximum.

Quantify it with illustrative (assumed, not hardware-measured) cycle costs: suppose if (threadIdx.x % 2 == 0) { pathA; } takes 20 cycles and the else branch takes 15 cycles. A non-divergent warp doing only pathA would finish in 20 cycles. A warp where half the lanes take each branch finishes in 20 + 15 = 35 cycles — a 35 / 20 = 1.75× slowdown, even though each individual lane did no more arithmetic than before. Notice, though, that the vectorAdd boundary check traced above does not cause this: with threadsPerBlock = 256 (eight warps of 32), the invalid region in the last block starts at threadIdx.x = 64, which is itself a multiple of 32 — so warps 0–1 are entirely valid and warps 2–7 are entirely invalid, and no single warp contains a mix. The boundary check is warp-uniform here purely because 64 happens to be warp-aligned; that alignment is not guaranteed in general, which is exactly what the active-recall question below tests.

Active recall

Attempt each question before reading its answer.

1. In the vectorAdd example (N = 1,000,000, threadsPerBlock = 256), what global index does the thread with threadIdx.x = 10 in blockIdx.x = 5 compute, and is it valid?
2. How many total threads does the launch vectorAdd<<<3907, 256>>> actually start, and how many of them never write to C?
3. Suppose threadsPerBlock changes from 256 to 128, with N still 1,000,000. Recompute blocksPerGrid, find the index of the last block, state how many of its threads are valid versus masked off, and say whether warp divergence occurs at that boundary.
4. A warp runs if (threadIdx.x % 2 == 0) { pathA (20 cycles) } else { pathB (15 cycles) }. How many cycles does the warp take in total, and what is the slowdown versus an idealized non-divergent warp running only pathA?
5. Why does shared memory exist at all, given that global memory can hold the same data and is visible to every thread?
6. A classmate says: "each CUDA thread runs independently, so an if/else inside a kernel never costs more than the branch itself." What is wrong with this claim?

Answers.

1. idx = blockIdx.x * blockDim.x + threadIdx.x = 5 * 256 + 10 = 1,290. Since 1,290 < 1,000,000, the thread is valid and writes C[1290] = A[1290] + B[1290].

2. Total launched threads: 3,907 × 256 = 1,000,192. Threads that never write: 1,000,192 - 1,000,000 = 192, all of them in the final block (blockIdx.x = 3906), with threadIdx.x from 64 to 255.

3. blocksPerGrid = ceil(1,000,000 / 128) = ceil(7,812.5) = 7,813. The last block is blockIdx.x = 7812, with offset 7,812 × 128 = 999,936 — the same offset as before, because 999,936 is a multiple of both 128 and 256. Valid threads: threadIdx.x = 0 to 63 (64 threads, reaching idx = 999,999); masked-off threads: threadIdx.x = 64 to 127 (64 threads). With blockDim.x = 128 there are four warps per block (128 / 32 = 4); the invalid region again starts exactly at threadIdx.x = 64, a multiple of 32, so warps 0–1 are fully valid and warps 2–3 are fully invalid — still no divergence, for the same reason as the 256-thread case. (A further ripple worth noticing: had threadsPerBlock instead been chosen as 100, then 1,000,000 mod 100 = 0 exactly, giving blocksPerGrid = 10,000 with zero masked-off threads anywhere — the boundary problem disappears entirely for block sizes that evenly divide N.)

4. The warp executes both paths serially because it is one instruction stream in lockstep: 20 + 15 = 35 cycles total. Slowdown versus the idealized 20-cycle case: 35 / 20 = 1.75×.

5. Shared memory is on-chip SRAM local to the SM running a given block, so its latency is roughly one to two orders of magnitude lower than global DRAM and it lets threads in the same block reuse data they have already fetched (for example, a tile of a matrix in tiled matrix multiplication) instead of every thread re-issuing its own slow global-memory read for the same value. The tradeoff is size and scope: it is tiny (up to 192 KB per SM on an A100, shared among all resident blocks on that SM) and invisible outside the block that allocated it, unlike global memory's full-device, multi-gigabyte reach.

6. The claim conflates a CUDA thread with an OS thread. CUDA schedules threads in fixed groups of 32 (a warp) that execute in SIMT lockstep — every active lane in the warp runs the same instruction on the same clock. An if/else that splits a warp's threads across both branches forces the hardware to execute both branches serially, masking off the inactive lanes each time, so the warp's total time is the sum of both branches' costs rather than the cost of one branch alone. Branch divergence is real and it is not free.

Think About It

Think about this: How would you explain cuda programming basics: gpu computing fundamentals to a friend who has never seen a computer? What real-world analogy would you use? Imagine you had to build a system using these concepts — what would be your first step? Try this: before moving on, write down three things you learned and one question you still have.

Practice Exercises

Now it is time to practice! Complete these challenges to solidify your understanding:

• Exercise 1: Write a short program that demonstrates the core concept from this chapter. Test it with at least 3 different inputs.
• Exercise 2: Find a real-world example where cuda programming basics: gpu computing fundamentals is used in an Indian company (like TCS, Infosys, Flipkart, or ISRO). Write a paragraph explaining the connection.
• Exercise 3: Create a mind-map connecting cuda programming basics: gpu computing fundamentals to at least 3 other topics you have studied.

Key Takeaways — Summary and Recap

Let us recap what we covered: the core ideas behind cuda programming basics: gpu computing fundamentals, how they connect to real-world applications, and why they matter for your journey in computer science. Remember these key points as you move forward. For competitive exam preparation (CBSE, JEE, BITSAT), focus on understanding the WHY behind each concept, not just the WHAT.