CUDA BASIC

This post is written for people who mainly use PyTorch but want to understand “what’s really happening” on the GPU.

https://github.com/Sung-Jae-Seong/CUDA_EXAMPLE

GPU = Heterogeneous Computing

CUDA programming is fundamentally a Host + Device model.

• Host (CPU + RAM) : Controls the overall program flow Input: transfers data to the GPU Launch: issues commands (kernel launches) to the GPU Output: retrieves results back from the GPU
• Device (GPU + VRAM)

Executes the simple, repetitive work in parallel (what the host “tells it to do”)

AI Hardware Spectrum (why GPUs matter)

AI hardware spectrum diagram

AI accelerators are dedicated hardware designed to speed up deep learning operations (especially massively-parallel math like matrix multiplication) with high throughput and efficiency.

• CPU: general-purpose; flexible; relatively weak at massive parallelism
• GPU: thousands of simpler cores; excellent for large-scale parallel workloads (DL training/inference throughput is high)
• Habana Gaudi: Intel’s training accelerator aimed at efficient large-scale distributed training
• AWS Inferentia: AWS inference chip optimized for low latency and high inference throughput
• FPGA: programmable hardware; flexible but harder to program
• ASIC: fixed-function; highest performance and power efficiency (but not reconfigurable)
• “General Purpose”: flexible for many workloads
• “Programmability”: how freely developers can modify/optimize behavior in software/hardware
• “Specialized”: designed for a narrow set of tasks, trading flexibility for speed/efficiency
• “Efficiency”: doing the same work with less power/resources

CPU vs GPU

A CPU is built for versatile, branch-intensive tasks using a small number of high-performance cores, deep cache hierarchies, low-latency memory access, and higher clock speeds, while a GPU is designed for large-scale parallel computation with many simpler cores, lower clock speeds, higher memory bandwidth, and dedicated VRAM, enabling it to handle massively parallel workloads such as graphics rendering and deep learning more efficiently. In certain AI servers equipped with multiple GPUs, the combined VRAM capacity can exceed the total system RAM.

• Typical Deeplearning Workflow

1. Data load: CPU loads from disk into RAM
2. Model definition: CPU constructs model; then .to(device) moves weights to VRAM
3. Data batch: batches are read from RAM and moved to VRAM
4. Forward pass, loss, backward, weight update: mostly executed on GPU
5. Convenience tasks (TensorBoard logging, checkpoint saving, etc.): typically on CPU

What is CUDA?

CUDA is NVIDIA’s parallel computing platform and programming model. It lets developers use GPUs for general-purpose computation (not just graphics), including deep learning, scientific computing, and HPC.

• Key CUDA terms
• Event: a synchronization marker for timing and stream coordination
• Stream: an ordered sequence of GPU operations (operations in a stream execute in issue order)
• CUDA Core: executes instructions for threads; a single CUDA core executes one thread’s instruction at a time (parallelism comes from having many cores + many threads)
• Streaming Multiprocessor (SM): the core execution unit on NVIDIA GPUs; schedules and runs warps
• Warp: a group of 32 threads that execute the same instruction at the same time (SIMT)
• malloc / cudaMalloc: memory allocation (cudaMalloc allocates device/global memory)
• global: marks a kernel launched by the host and executed on the device
• restrict: tells the compiler the pointer is the only reference to that object (enables stronger optimizations like vectorization)
• blockDim: threads per block (a thread block can have up to 1024 threads)
• Grid: a collection of blocks launched for one kernel call (in Triton, “grid” often refers to the number of parallel program instances)

SM / Warp / Grid / Block / Thread

SM

• An NVIDIA GPU is built from many SMs. Each SM includes multiple execution resources such as: FP32 units: 32-bit floating-point arithmetic INT32 units: integer arithmetic (addressing, indexing, etc.) Tensor Cores: matrix-multiply acceleration units (deep learning) SFU: special functions (sqrt, trig, etc.) Register file: per-thread register storage RT core: ray-tracing acceleration (graphics) Shared memory / L1 cache: fast on-SM memory shared by threads in a block

Note: many NVIDIA architectures can execute fused multiply-add (a * b + c) as a single instruction; throughput depends on the specific GPU generation.

Warp

warp scheduling diagram

Threads are not scheduled one-by-one. They are typically executed in groups of 32 threads called a warp.

Threads in the same warp run the same instruction simultaneously. Warp scheduler chooses which warp to run next. Instruction Dispatch Unit sends the selected warp’s instruction to execution units. Different warps interleave execution: while one warp waits on memory, another warp can run, reducing idle time.

Grid / Block / Thread

When the host launches a kernel, one Grid is created per launch.

• Grid = collection of Blocks
• Block (thread block) = collection of Threads
• Thread = the smallest execution entity

kernel launch -> grid creation

Each kernel launch can use a different grid size (different block counts). Threads have private registers. Each block has shared memory (shared among threads in the block). “Global memory” (device memory from cudaMalloc) is accessible by all threads across the grid. Host launches kernel → Grid is created → Grid contains Blocks → Blocks are assigned to SMs → SM runs Threads in Warps.

• Physical vs logical SM = physical hardware Grid / Block / Thread / Warp = logical execution hierarchy

CUDA Concurrency

Concurrency Example

Amount of Concurrency

If you do Host-to-Device copy (H2D), then run the kernel, then Device-to-Host copy (D2H) strictly in sequence, there can be a lot of wasted time (CPU or GPU sitting idle).

H2D: host (CPU) → device (GPU) / D2H: device (GPU) → host (CPU)

With CUDA streams and asynchronous memcpy, you can overlap:

• H2D transfers
• Kernel execution
• D2H transfers
• Plus unrelated CPU work

People often describe this as “2-way concurrency” (copy + compute overlap), “3-way concurrency” (H2D + compute + D2H overlap), etc. In practice, speedups depend on hardware engines, transfer sizes, kernel shape, and whether your pipeline can actually be overlapped.

Profiling

When you launch a CUDA kernel, the CPU does not “wait” for it by default. The CPU issues the launch and continues immediately. That means Python-level timers (like time.time) can easily measure only the launch overhead, not the real GPU execution time.

Correct timing approaches are

• Force synchronization (torch.cuda.synchronize) around the measured region.
• Use torch.cuda.Event, which is designed to measure device time.

Profiling lets you see how work splits between CPU and CUDA, and helps identify:

• Host-to-device copies
• Kernel launch overhead
• Which ops dominate runtime

A common best practice is to “skip” early iterations and only record the meaningful steady-state region:

• wait: ignore first N steps
• warmup: prepare without recording
• active: record real traces
• repeat: repeat that cycle
• Also the first CUDA kernel run is often slow due to driver loading, allocations, and JIT preparation.

Experiment

Running a custom CUDA kernel from PyTorch (load_inline) it makes the mechanics explicit.

1. Write C++/CUDA code as a string inside Python
2. JIT compiles it with nvcc and builds a shared library
3. PyBind connects it as a Python module
4. You call it like a regular Python function

This is the minimum you need to see “my code is actually running on the GPU,” not just calling existing PyTorch ops.

Indexing math you should memorize

• Almost every beginner CUDA kernel starts from the same indexing rule:
• Global ID = blockIdx.x × blockDim.x + threadIdx.x
• This single integer decides which element (or pixel) each thread owns. Also, boundary checks matter. Your grid rarely matches the data length perfectly, so you typically guard with code like: if i < n.

Example code : RGB 2 GRAY

import torch
from torch.utils.cpp_extension import load_inline  # JIT compile and load C++/CUDA extension at runtime
import matplotlib.pyplot as plt
cuda_source = \
r'''
#include <torch/extension.h>      // PyTorch C++ API
#include <cuda.h>                 // CUDA driver API
#include <cuda_runtime.h>         // CUDA runtime API
__global__ void rgb_to_gray_kernel(const float* input, float* out, int n) {
    // __global__ : kernel function executed on GPU, launched from host
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // blockIdx.x : index of current block in the grid (1D grid here)
if (i < n) {
        float r = input[i * 3 + 0];
        float g = input[i * 3 + 1];
        float b = input[i * 3 + 2];
        out[i] = 0.21f * r + 0.72f * g + 0.07f * b;
    }
}
torch::Tensor rgb_to_gray(torch::Tensor input) {
    auto H = (int)input.size(0);
    auto W = (int)input.size(1);
    int n = H * W;
    // Allocate GPU output tensor using same device/options
    auto output = torch::empty({H, W}, input.options().dtype(torch::kFloat32));
    const int threads = 256;                  // Threads per block
    const int blocks = (n + threads - 1) / threads;  // Grid size
    // CUDA kernel launch: <<<gridDim, blockDim>>>
    rgb_to_gray_kernel<<<blocks, threads>>>(
        (const float*)input.data_ptr<float>(),   // Device pointer access
        (float*)output.data_ptr<float>(),
        n
    );
    return output;
}
'''
cpp_header = \
r'''
torch::Tensor rgb_to_gray(torch::Tensor input);
'''
module = load_inline(
    name='rgb_to_gray_cpp',
    cpp_sources=cpp_header,
    cuda_sources=cuda_source,
    functions=['rgb_to_gray'],
    verbose=False,
)
# load_inline : compiles and dynamically loads CUDA/C++ code as a PyTorch extension
def rgb_to_gray(img_hwc):
    return module.rgb_to_gray(img_hwc)
if __name__ == "__main__":
    device = "cuda"  # Execute tensor operations on GPU
    H, W = 1080, 1920
    img = torch.zeros((H, W, 3), device=device, dtype=torch.float32)
    w3 = W // 3
    img[:, :w3, 0] = 255.0
    img[:, w3:2*w3, 1] = 255.0
    img[:, 2*w3:, 2] = 255.0
    img = img.contiguous()  # Ensure contiguous memory layout for raw pointer access
    gray = rgb_to_gray(img)
    torch.cuda.synchronize()  # Synchronize CPU with GPU execution
    print(img.shape, gray.shape, gray.dtype)
    img_cpu = img.byte().cpu().numpy()   # Device → host transfer
    gray_cpu = gray.cpu().numpy()
    plt.figure(figsize=(10, 4))
    plt.subplot(1, 2, 1)
    plt.title("RGB")
    plt.imshow(img_cpu)
    plt.axis("off")
    plt.subplot(1, 2, 2)
    plt.title("Grayscale")
    plt.imshow(gray_cpu, cmap="gray")
    plt.axis("off")
    plt.tight_layout()
    plt.show()