A piece of GPU code

__global__

In the main code, the function is called like:

vec_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

This function runs on the GPU, executes many threads at the same time. Each thread runs the same code, but with difference IDs. This is SIMT: singe instruction, multiple threads.

1️⃣ What is a kernel (in CUDA)?

In CUDA:

A kernel is a function that runs on the GPU and is executed by many threads in parallel.

That’s it.

Formally:

• A kernel is a function declared with __global__
• It is launched from the CPU
• It runs on the GPU
• It is executed once per GPU thread

<<<>>> is syntax of lunching CUDA kernel.

Understanding:

Each kernel thread already has an identity (blockIdx, threadIdx).
The variable i is a user-defined mapping from that identity to a unit of work (e.g., an array index).

✅ “each parallel thread computes its own local variable i”

• The kernel is launched once
• The kernel code runs in many threads
• Each thread has its own registers
• i lives in a register, private to that thread

There is no shared global i.

What i actually is (hardware perspective)

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

This line:

• Executes independently in every thread
• Uses hardware-provided registers:
  • blockIdx
  • threadIdx
• Produces a per-thread scalar
• Stored in a register (not memory)

So yes:

i is a logical work ID, not a hardware ID.

The general form (memorize this)

kernel_name<<<grid_dim, block_dim>>>(arguments);

More fully:

kernel_name<<<num_blocks, threads_per_block>>>(args...);

CUDA organizes threads in a 3-level hierarchy:

Grid
└── Block
└── Thread

__global__, __device__, and __host__ are language-level qualifiers, not function names.
You must use the exact ones CUDA defines.

Thread, Block and Grid

Thread (smallest unit)

• Each thread has: its own registers, its own local variables, Threads are very lightweight

🧱 Block (group of threads)

• A block is a group of threads
• Threads in the same block can share fast memory, can synchronize

🌍 Grid (all blocks)

• A grid is all blocks launched for a kernel, This is the entire kernel execution

Total threads = blocks × threads

The built-in CUDA variables

These are special variables provided by CUDA.
You do not define them — the GPU runtime does.

threadIdx

blockDim

• Number of threads per block
• Same value for all threads
• Chosen by you at launch

vec_add<<<blocks, 256>>>(...)

blockIdx

• Block’s index inside the grid
• Range: 0 → gridDim.x - 1

How this maps to real GPU hardware

This is the architecture part you asked for 👇

GPU cores are organized into SMs (Streaming Multiprocessors)

Think:

• CPU: few powerful cores
• GPU: many SMs, each running many threads

Each SM:

• Executes one or more blocks
• Schedules threads in groups of 32 threads called a warp

Warp (very important)

• A warp = 32 threads
• Warps execute the same instructionon different data

Why block size matters

Good block sizes:

• multiples of 32
• common choices:
  • 128
  • 256
  • 512

This keeps hardware fully utilized.

Why the double underscore syntax exists

CUDA chose __keyword__ to:

• Avoid clashes with C++ keywords
• Make CUDA extensions obvious

Other compilers will just ignore or reject them.

Mental model (important)

Rule of thumb (very useful)

Only kernels use __global__.
Everything else on GPU is __device__.

vec_add code example on GPU:

#include <cstdio>
#include <cuda_runtime.h>

__global__ void vec_add(const float* a, const 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() {
    int n = 1 << 20;
    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] = i;
        h_b[i] = 2 * i;
    }

    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;
    vec_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    printf("h_c[123] = %f\n", h_c[123]);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
}