Technology 📅 Sep 07, 2026 👁️ 324 views

CUDA Explained: The Programming Model Behind GPU-Accelerated AI

A
Admin User

admin

What CUDA actually is, how the execution and memory models work under the hood, and how the same idea shows up whether you're writing raw kernels in C++ or calling a one-line NumPy-style function in Python.
2007
CUDA first released
32
Threads per warp, always
4
Abstraction levels covered below

If you've worked with PyTorch or TensorFlow, you've already used CUDA — you just haven't seen it directly. This post pulls back that curtain: what CUDA is, how NVIDIA GPUs actually organize parallel work, and how the same underlying model shows up whether you're writing a raw kernel in C++ or letting CuPy handle everything with a single line of Python.

What Is CUDA?

CUDA (Compute Unified Device Architecture) is a parallel computing platform and API created by NVIDIA. It lets developers use CUDA-enabled GPUs for general-purpose computing — not just graphics — a category of work usually called GPGPU (general-purpose computing on GPUs).

Practically, CUDA is what lets you write code that runs on thousands of GPU cores simultaneously instead of writing a normal sequential program for a CPU. Every major deep learning framework — PyTorch, TensorFlow, JAX — is built on top of it, usually through NVIDIA's Deep Learning SDK, which includes libraries like cuDNN for the neural network primitives those frameworks rely on internally.

CUDA only runs on NVIDIA GPUs. This is the ecosystem lock-in effect covered in Part 1 of this series — once frameworks defaulted to CUDA, the entire deep learning tooling stack became built around NVIDIA hardware specifically.
The CUDA Toolkit and Version Management

CUDA itself is installed via the CUDA Toolkit, and this is where a lot of developers hit their first real friction point: CUDA versions matter, a lot. Specific hardware supports specific ranges of CUDA versions, and specific libraries or frameworks often only work correctly with specific CUDA versions. It's genuinely common to need multiple CUDA versions installed side by side on the same machine.

Here's the catch: CUDA doesn't ship its own version manager. Instead, the standard practice is environment isolation using Conda, Mamba, or Docker.

Creating an isolated environment with a specific CUDA version
conda create -n cuda118 python=3.10 cudatoolkit=11.8
conda activate cuda118
Or isolate entirely with Docker
docker run --gpus all -it nvidia/cuda:11.8.0-devel-ubuntu22.04

Docker in particular has become the default for anything production-bound, since it guarantees the exact CUDA version, driver compatibility, and library set travels with the container rather than depending on whatever happens to be installed on the host machine.

Understanding the Execution Hierarchy

This is the part of CUDA that trips up most people coming from regular CPU programming, so it's worth being precise about the terminology.

The CUDA Execution Hierarchy GRID all blocks launched by one kernel call BLOCK a group of threads scheduled together WARP 32 threads executing in lockstep (THREADS) A block can contain many warps — this is one of several in the block. Threads in a block can share memory and synchronize. BLOCK, BLOCK, BLOCK… The CUDA runtime schedules blocks onto whichever Streaming Multiprocessors (SMs) are free — you don't control the mapping. Blocks CANNOT share memory or synchronize with each other. More SMs in the GPU = more blocks can run concurrently.
Term
What It Means
Thread
The lowest-level unit of parallel execution — one single instance of your kernel function
Warp
A fixed group of 32 threads that execute the same instruction at the same time on a Streaming Multiprocessor
Block (Thread Block)
A group of threads that execute together and can share memory and synchronize with each other
Grid
A group of thread blocks — a single grid maps to one GPU, and can span multiple SMs
SM (Streaming Multiprocessor)
The core execution unit inside the GPU that actually runs warps of threads; contains CUDA cores, Tensor Cores, and load/store units

The relationship that matters most: you write code for a single thread, and CUDA runs that same code across every thread in a block, and every block in a grid. The CUDA runtime decides how to distribute blocks across the GPU's available SMs — you don't control that mapping directly, and it can vary based on how many SMs the specific GPU has.

A CUDA core is not the same thing as a CPU core. It's a much smaller, simpler arithmetic unit, handling basic floating-point and integer math (FP32, INT32) for the threads assigned to it — not a general-purpose processor in its own right.
CUDA Memory Hierarchy

Just like execution, CUDA memory is organized in a hierarchy — and picking the right memory tier for the right job is one of the biggest levers for performance.

CUDA Memory Hierarchy: Speed vs. Size Trade-off Registers Shared / L1 L2 Cache Global Memory (GPU DRAM) Fastest ~1 cycle Slowest ~400–800 cycles Smallest per-thread Largest GBs, shared by all threads Registers — private per thread, compiler-managed Shared Memory / L1 — shared within a block, ~48–96KB per block L2 Cache — shared across all SMs, 40MB on A100 Global Memory — accessible by all threads, persists for the program
Memory Type
Scope
Speed
Typical Size
Registers
Private to each thread
~1 cycle (fastest)
Very small, per-thread
Shared Memory / L1
Shared within a block
A few cycles
~48–96KB per block
L2 Cache
Shared across all SMs
Fast
40MB on A100, 6MB on V100
Global Memory
Accessible by all threads, persists for the program
~400–800 cycles (slowest)
GBs — the GPU's full DRAM

There's also constant memory (read-only, cached, very fast if every thread reads the same value, limited to 64KB) and local memory, which is private per-thread but actually physically lives in global memory — used automatically when a thread's registers overflow. That register spillover is a common, often invisible, source of slowdowns in poorly tuned kernels.

The CUDA Processing Flow

Every CUDA program, regardless of what it's doing, follows the same basic four-step pattern:

The Standard CUDA Processing Flow 1. Copy data: Host → Device CPU memory to GPU memory over PCIe 2. CPU launches the kernel Specifies grid & block dimensions 3. CUDA cores execute in parallel Thousands of threads run the same kernel code 4. Copy result: Device → Host GPU memory back to CPU memory Steps 1 and 4 (the PCIe transfers) are often the real bottleneck — not the compute itself. This is why GPUDirect and unified memory exist.

This flow explains why naive GPU code can sometimes be slower than CPU code for small workloads — the PCIe transfer overhead in steps 1 and 4 can dominate if the actual computation in step 3 is trivial. It's also the reason technologies like GPUDirect RDMA and GPUDirect Storage exist: they let data move directly between GPU memory and network or storage devices, bypassing the CPU entirely and cutting out one of these round trips.

Debugging and Profiling Tools

When a kernel misbehaves or underperforms, these are the standard tools for figuring out why:

Tool
What It's For
NVIDIA Nsight
Low-overhead profiling, tracing, and debugging — the primary tool for understanding where time is actually going in a kernel
CUDA-GDB
An extension of the standard Linux GDB debugger, giving you a console-based debugging interface for CUDA code specifically
CUDA Memcheck
Surfaces memory access issues — out-of-bounds reads/writes, misaligned access — across the thousands of threads running concurrently, which is nearly impossible to catch by inspection alone
Writing CUDA Code: Four Levels of Abstraction

You don't have to write raw CUDA C++ to use CUDA. There's a whole ladder of abstraction, and where you land on it is mostly a trade-off between control and productivity. Here's the exact same vector-addition operation, written four different ways.

1. Raw CUDA C++

Full manual control: you write the kernel, manage device memory allocation, and handle host-to-device transfers explicitly.

vector_add.cu
__global__ void vectorAdd(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];
    }
}

int main() {
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
}

2. PyCUDA

The kernel itself is still raw CUDA C, but it's embedded as a string inside Python, and PyCUDA handles a lot of the boilerplate around memory management.

vector_add_pycuda.py
import pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
import numpy as np

mod = SourceModule("""
__global__ void vectorAdd(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];
}
""")

vectorAdd = mod.get_function("vectorAdd")
vectorAdd(a_gpu, b_gpu, c_gpu, np.int32(n),
          block=(256, 1, 1), grid=(blocks, 1, 1))

3. Numba

The kernel is now actual Python, compiled Just-In-Time (JIT) to run on the GPU via the @cuda.jit decorator. No CUDA C required at all — and memory management becomes largely automatic.

vector_add_numba.py
from numba import cuda
import numpy as np

@cuda.jit
def vector_add(a, b, c):
    i = cuda.grid(1)
    if i < a.size:
        c[i] = a[i] + b[i]

vector_add[blocks_per_grid, threads_per_block](a, b, c)

4. CuPy

No kernel, no explicit memory management — CuPy mirrors the NumPy API almost exactly, and operations like addition just run on the GPU transparently. This is the level most data scientists actually work at day to day.

vector_add_cupy.py
import cupy as cp

a = cp.array([1, 2, 3, 4, 5])
b = cp.array([10, 20, 30, 40, 50])
c = a + b  # runs on the GPU, no kernel code needed
Notice the pattern: as you move up this ladder, you trade fine-grained control for productivity. Raw CUDA C++ gives you the most performance headroom for a hand-tuned kernel; CuPy gets you 90% of the benefit with a fraction of the code. Most real projects end up using a mix — CuPy or Numba for most operations, with a hand-written kernel only for the one hot loop that actually needs it.
The CUDA Library Ecosystem

On top of the core CUDA runtime, NVIDIA ships a large set of pre-built libraries so you rarely need to write low-level kernels for common operations:

Library
What It Does
cuBLAS
CUDA Basic Linear Algebra Subroutines — matrix multiplication and friends
cuFFT
Fast Fourier Transform on the GPU
cuRAND
GPU-accelerated random number generation
cuSOLVER
Dense and sparse direct solvers
cuSPARSE
Operations on sparse matrices
NPP
NVIDIA Performance Primitives — image and signal processing
NVML
NVIDIA Management Library — the library behind tools like nvidia-smi
cuDNN
Deep Neural Network primitives — convolution, pooling, normalization, activation functions. This is what PyTorch and TensorFlow call under the hood.

cuDNN deserves a special mention: it's technically separate from the core CUDA Toolkit, distributed as part of NVIDIA's Deep Learning SDK, and it's the single library most responsible for how fast your PyTorch or TensorFlow training loop actually runs. When people talk about a GPU being "well optimized for deep learning," a large part of what they mean is that cuDNN has tuned, hardware-specific implementations of the exact operations those frameworks use most.

Recap: what to take away

  • CUDA lets GPUs run general-purpose parallel code, not just graphics — and it's the foundation every major deep learning framework is built on.
  • The execution hierarchy (thread → warp → block → grid) is really about how the same kernel code gets replicated across thousands of parallel threads, scheduled by the GPU's SMs.
  • Memory tier choice matters: registers and shared memory are fast but tiny; global memory is huge but slow — performance tuning is mostly about minimizing trips to global memory.
  • You don't have to write raw CUDA C++. PyCUDA, Numba, and CuPy each trade some control for a lot of productivity, and most real projects mix levels depending on the task.
PART 3 OF 5 · NVIDIA SERIES Next: Building an AI Data Center — H100, DGX Systems, and the Power/Cooling Reality →
Recently Enrolled

Student enrolled in this course.

View course
Explore Courses

Latest from @kp__expert

Follow on Instagram
Loading Instagram posts...

AI Course Assistant

Share your details and goals to get the best course recommendations.

Recommended Courses

Select a course name to view full details.

Course Details
Enrollment & Contact
  • Review selected course and confirm your enrollment request.
  • Click checkout to move into the full payment process.
  • After payment submission, your enrollment is processed by our team.
Admissions Contact
Email: info@kpexpert.com
Phone: +91 92708 37105
Your submitted details
Name, email and phone will appear here.
Your request has been submitted successfully. Our team will contact you shortly.