From Threads to Blocks: Fundamental CUDA Concepts Explained#

Threads: The Smallest Unit#

A thread is the smallest unit of execution on a GPU. Each thread runs a particular instance of a kernel. Compared to a CPU thread, GPU threads are more lightweight, and you can have thousands or millions of them active at once.

Warps: A Hardware Concept#

When you request a certain number of threads, the GPU hardware will schedule them in groups called warps (typically of size 32 threads on NVIDIA GPUs). All threads in a warp execute the same instruction simultaneously (SIMT, or Single Instruction, Multiple Threads). Divergence within a warp (e.g., divergent if statements) can reduce efficiency.

Why Blocks Matter#

Threads are grouped into blocks. A block is an array (1D, 2D, or 3D) of threads, and it provides:

• A rich set of thread indexing capabilities.
• Shared memory for better data sharing among threads in the same block.
• Synchronization mechanisms such as __syncthreads().

Thread Indexing Within a Block#

Each thread within a block has an ID accessible via CUDA built-in variables like:

• threadIdx.x, threadIdx.y, threadIdx.z (the thread’s coordinate within the block).
• blockDim.x, blockDim.y, blockDim.z (the block’s size along each dimension).

Typically, you compute a global index when accessing data in memory:

1
__global__ void myKernel(float *data) {
2
    // Compute the global thread index for a 1D grid
3
    int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x;
4
    float value = data[globalThreadId];
5
    // ... do something with value
6
}

In this example, globalThreadId uniquely identifies each thread across the entire grid. For multi-dimensional scenarios, you extend this logic with blockIdx.y, threadIdx.y, and so on.

Block Size Considerations#

Selecting the right block size is crucial for performance. Key tips:

• Typically, you want each block to have a number of threads that is a multiple of the warp size (32).
• Common block sizes include 128, 256, or 512 threads per block.
• The maximum number of threads per block is GPU-dependent (up to 1024 on many modern GPUs).

Grid Dimensions#

• gridDim.x, gridDim.y, and gridDim.z store the number of blocks along each dimension.
• blockIdx.x, blockIdx.y, and blockIdx.z identify the block’s position within the grid.

Conceptually:

1. You define how many blocks you want in your grid.
2. You define how many threads go in each block.

For instance, suppose you have 1024 elements to process and decide to run 256 threads per block. That means you need 4 blocks in total for a 1D arrangement:

1
dim3 blocks(4);
2
dim3 threads(256);
3
myKernel<<<blocks, threads>>>(deviceData);

Here, gridDim.x = 4, blockDim.x = 256, and total threads = 4 * 256 = 1024.

Global Memory#

The largest and slowest memory space. Kernel arguments and large arrays often reside here.

Shared Memory#

A fast, on-chip memory shared by threads within the same block. Proper usage can significantly improve performance, but it’s limited in size (commonly tens of kilobytes per block).

Registers#

Each thread has access to a set of registers. They are extremely fast but limited. Overusing registers might spill data into local memory, which is stored in global memory.

The Kernel#

1
__global__ void addVectors(const float *a, const float *b, float *c, int n) {
2
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
3
    if (idx < n) {
4
        c[idx] = a[idx] + b[idx];
5
    }
6
}

• We calculate idx based on the block and thread indices.
• We check if (idx < n) to avoid out-of-bounds memory access.

The Host Code#

1
#include <iostream>
2
#include <cuda.h>
3

4
__global__ void addVectors(const float *a, const float *b, float *c, int n);
5

6
int main() {
7
    int n = 1 << 20;  // 1 million elements
8
    size_t size = n * sizeof(float);
9

10
    // Allocate host memory
11
    float *h_a = (float*)malloc(size);
12
    float *h_b = (float*)malloc(size);
13
    float *h_c = (float*)malloc(size);
14

15
    // Initialize host arrays
16
    for(int i = 0; i < n; i++) {
17
        h_a[i] = 1.0f;
18
        h_b[i] = 2.0f;
19
    }
20

21
    // Allocate device memory
22
    float *d_a, *d_b, *d_c;
23
    cudaMalloc((void**)&d_a, size);
24
    cudaMalloc((void**)&d_b, size);
25
    cudaMalloc((void**)&d_c, size);
26

27
    // Copy data to device
28
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
29
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
30

31
    // Define block and grid dimensions
32
    int blockSize = 256;
33
    int gridSize = (n + blockSize - 1) / blockSize;
34

35
    // Launch kernel
36
    addVectors<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
37

38
    // Copy result back to host
39
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
40

41
    // Verify
42
    for(int i = 0; i < 10; i++) {
43
        std::cout << h_c[i] << " ";
44
    }
45
    std::cout << std::endl;
46

47
    // Cleanup
48
    free(h_a); free(h_b); free(h_c);
49
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
50
    return 0;
51
}
52

53
__global__ void addVectors(const float *a, const float *b, float *c, int n) {
54
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
55
    if (idx < n) {
56
        c[idx] = a[idx] + b[idx];
57
    }
58
}

Declaring Shared Memory#

Within a __global__ or __device__ function, you can declare a shared memory array:

1
__global__ void kernelWithShared(float *data) {
2
    __shared__ float tile[256];  // This is allocated per block
3
    int idx = threadIdx.x;
4
    tile[idx] = data[idx];
5
    __syncthreads();
6

7
    // Now all threads in this block can read tile[]
8
    float val = tile[(idx+1) % 256];
9
    // ...
10
}

• __syncthreads() is crucial to ensure all writes to shared memory are visible to all threads in the block.

When to Use Shared Memory#

• The data must be reused multiple times within a block.
• The size of data is within the hardware’s shared memory limit (often 48KB or 96KB per multiprocessor, depending on configuration and GPU generation).

Thread Synchronization#

CUDA provides multiple mechanisms for synchronization:

1. __syncthreads() - Ensures all threads in a block reach this point before continuing.
2. __syncwarp() - Synchronizes threads in a warp (on GPUs with compute capability >= 7.0, you can specify a mask).

Going beyond block-level synchronization typically requires splitting the operations into multiple kernels or using more advanced concurrency features.

Atomic Operations#

When multiple threads need to update shared data concurrently, you can use atomic operations:

1
__global__ void atomicAddKernel(int *array) {
2
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
3
    atomicAdd(&array[0], idx);
4
}

CUDA’s atomicAdd, atomicSub, atomicMax, etc., ensure data integrity, but may reduce performance if contention is high.

What is a Stream?#

A stream is a sequence of operations (kernels, memory copies, etc.) that execute in order on the GPU. By default, kernels run in stream 0, which is synchronous with respect to host code in many cases.

Overlapping Operations#

Using multiple streams allows for concurrent:

• Kernel execution in one stream,
• Memory copy in another stream,
• Or simply different kernels executing in parallel if resources permit.

1
cudaStream_t s1, s2;
2
cudaStreamCreate(&s1);
3
cudaStreamCreate(&s2);
4

5
// Launch kernels in different streams
6
kernelA<<<grid, block, 0, s1>>>(...);
7
kernelB<<<grid, block, 0, s2>>>(...);
8

9
// Non-blocking if events or synchronization are not used
10
cudaMemcpyAsync(..., s1);
11
cudaMemcpyAsync(..., s2);
12

13
// Cleanup
14
cudaStreamDestroy(s1);
15
cudaStreamDestroy(s2);

Effective use of streams can drive better GPU occupancy and overall throughput.

Constant Memory#

Constant memory is cached read-only memory. If many threads read the same value from constant memory, the caching mechanism can reduce global memory bandwidth usage. You declare it like:

1
__constant__ float constData[256];

And copy from host to device with:

1
cudaMemcpyToSymbol(constData, hostData, size);

Texture Memory#

Texture memory is specialized and also cached, often used for 2D and 3D data. It provides built-in filtering and addressing modes. While it’s historically associated with graphics, it can also boost performance for certain data access patterns in GPGPU workloads.

Dynamic Parallelism#

With dynamic parallelism, kernels can launch other kernels directly from the GPU. For example:

1
__global__ void childKernel() {
2
    // ...
3
}
4

5
__global__ void parentKernel() {
6
    // Launch child kernel from within the GPU
7
    childKernel<<<1, 32>>>();
8
}
9

10
int main() {
11
    // Launch the parent kernel
12
    parentKernel<<<1, 1>>>();
13
    cudaDeviceSynchronize();
14
    return 0;
15
}

This feature can simplify complex workflows where parallel work spawns more parallel work. However, it can also introduce overhead and complicate resource management.

Unified Memory#

Unified memory automatically manages data across CPU and GPU, introduced with CUDA 6 and higher. It simplifies memory handling:

1
float *unifiedData;
2
cudaMallocManaged(&unifiedData, n * sizeof(float));
3
// Access from both host and device without explicit cudaMemcpy

But for performance-critical applications, manual memory management may yield better results.

Occupancy and Resource Considerations#

• Occupancy refers to how many warps can run concurrently on a streaming multiprocessor (SM).
• You can tune thread block sizes, shared memory usage, and registers to improve occupancy.

Coalesced Global Memory Access#

Optimize global memory accesses such that consecutive threads access consecutive memory addresses. This is called coalescing and drastically improves bandwidth utilization.

Profiling Tools#

NVIDIA provides several profiling and analysis tools:

• NVIDIA Nsight Compute: A low-level kernel profiler.
• NVIDIA Nsight Systems: A system-wide profiler to see how CPU and GPU tasks are scheduled over time.

Use these tools to find bottlenecks in memory bandwidth, compute, or other areas.

Further Reading and References#

• NVIDIA’s official CUDA Toolkit Documentation
• Mark Harris’s blog posts on Parallel Forall
• “Programming Massively Parallel Processors: A Hands-on Approach” by David B. Kirk and Wen-mei W. Hwu

Best of luck, and happy coding in CUDA!