Harnessing Horsepower: GPU Schedulers and Memory Hierarchies Explained#

Basic Concept of GPU Scheduling#

The core principle of GPU scheduling is that you have a massive set of threads, grouped in certain ways (warps, thread blocks, grids), and a hardware scheduler that picks which groups of threads get to execute on the GPU’s streaming multiprocessors (SMs). In most GPU programming models (e.g., CUDA, OpenCL), a kernel is launched with a specific configuration:

• Number of thread blocks (also known as work-groups in OpenCL).
• Number of threads within each block.

When a kernel is launched, the GPU hardware will dynamically schedule these blocks across the available compute units (SMs). Each SM in turn keeps track of the threads assigned to it and runs them in small groups known as warps (in CUDA) or wavefronts (in AMD terminology).

Warp Schedulers#

A warp is a bundle of 32 threads (for NVIDIA GPUs) or 64 threads (for AMD GPUs) that execute in lockstep. Each instruction cycle, the scheduler issues instructions to one warp. However, if the warp is waiting on data from memory, the scheduler can switch instantly to another warp that is ready to execute. This technique is called thread-level parallelism (TLP).

Key characteristics:

• SIMT (Single Instruction, Multiple Threads) model for NVIDIA GPUs.
• In one clock cycle, all threads in a warp are expected to execute the same instruction.
• Divergence (e.g., if some threads take a different branch) leads to serialization of instructions, which can degrade performance.

Warp schedulers manage which warp is active at any given cycle. They ensure maximum occupancy by hiding latency: if Warp A is waiting on memory, Warp B is chosen if it is ready. As a developer, it’s important to keep warps fully utilized and minimize thread divergence.

Thread Block Schedulers#

Above the warp level, the GPU organizes threads into thread blocks (or blocks, for short). Each block potentially contains multiple warps. The hardware tries to fill each Streaming Multiprocessor (SM) with as many blocks as possible until resources (registers, shared memory, etc.) are exhausted. The block-level scheduler ensures:

1. Resource Allocation: Each block receives a portion of register and shared memory resources.
2. Fairness: No single kernel hogs all the SMs if multiple kernels are being run in parallel.
3. Prioritization: Some GPUs allow priorities to be set, meaning high-priority tasks can preempt or out-prioritize lower-level tasks.

Modern GPUs allow overlapping of kernels and concurrency in various forms, making block scheduling dynamic.

Multi-Queue and Multi-Process Service (MPS)#

Modern data centers often require multiple processes to share a single GPU. This is where features like Multi-Process Service (MPS) in NVIDIA or Multi-Queue in AMD come into play. These techniques enable:

• Multiple CPU processes to share a single GPU context.
• Consolidation of multiple kernels from different processes into a single hardware queue.
• Improved utilization by filling the GPU with work from multiple clients.

MPS reduces overhead by allowing multiple applications to share GPU resources without extensive context switching. Shared contexts help keep the GPU busy, minimizing idle time. However, it requires careful resource allocation and consideration of how memory usage might collide among different processes.

Global Memory#

Global Memory is the largest pool of memory on a GPU (often several GBs to tens of GBs in modern GPUs). It corresponds to dedicated VRAM (Video RAM). Characterized by:

• High capacity (relative to other GPU memory levels).
• High latency (hundreds of clock cycles).
• Accessible by all threads across all SMs.

When a kernel reads or writes to global memory, memory transactions are coalesced whenever possible. Coalescing means that consecutive threads reading/writing consecutive addresses can form a single transaction, which is more efficient than issuing one transaction per thread.

Performance Tips:

• Keep data in global memory only when it needs to persist.
• Minimize random accesses to global memory.
• Aim for coalesced accesses to reduce transaction overhead.

Shared Memory#

Shared Memory resides on-chip within each Streaming Multiprocessor. It functions almost like a software-managed cache:

• Much lower latency than global memory.
• Limited in size (e.g., 48 KB, 64 KB, or 96 KB depending on GPU architecture and configuration).
• Accessible by all threads in a single thread block.

Threads within a block can collaborate by reading and writing shared memory. This makes algorithms like tiling or chunk-based operations extremely efficient. By chunking data into shared memory, threads can reuse data without repeatedly going to global memory.

For example, a matrix multiplication kernel might load sub-blocks (tiles) of both input matrices into shared memory, perform calculations on those tiles, and then write partial results back to global memory. This approach drastically reduces global memory traffic.

Registers and Local Memory#

Each thread has its own set of registers, which offer the lowest latency. Key points:

• Register usage is heavily impacted by compiler optimization.
• If a kernel uses many registers, fewer concurrent threads (warps) can be scheduled, reducing occupancy.

When a thread requires more storage than is available in registers, the GPU spills data into local memory (not to be confused with local methodology in CPU terms). In GPU terms, local memory resides in global memory space, meaning it has high latency. Thus, one typically aims to minimize register spilling to local memory whenever possible.

L1 and L2 Caches#

Modern GPUs also have hardware-managed caches:

• L1 Cache can be specialized or unified with shared memory, depending on configuration and GPU architecture (e.g., the “unified memory architecture” on some GPUs merges L1 with shared).
• L2 Cache is bigger than L1 and is shared across all SMs on the GPU. It helps reduce the global memory traffic for repeated accesses.

Because caching is hardware-managed, developers have less control over how data is cached, but they can optimize data layouts and access patterns to improve cache hit rates.

Constant and Texture Memory#

Two specialized memory types also exist on many GPU architectures:

1. Constant Memory: Read-only, often cached across multiple SMs, used for small, unchanging data.
2. Texture Memory: Also read-only in many cases, but optimized for certain access patterns (e.g., 2D locality). Frequently used in image processing or matrix manipulations.

These memories can significantly reduce bandwidth costs if your dataset can be stored in them and if you primarily read from them.

A Basic CUDA Kernel#

Below is an example of a simple CUDA kernel that doubles each element of an array:

1
#include <stdio.h>
2

3
__global__ void doubleArray(float* d_in, float* d_out, int size) {
4
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
5
    if (idx < size) {
6
        d_out[idx] = 2.0f * d_in[idx];
7
    }
8
}
9

10
int main() {
11
    int size = 1 << 20; // 1 million elements
12
    size_t bytes = size * sizeof(float);
13

14
    // Allocate host memory
15
    float* h_in = (float*)malloc(bytes);
16
    float* h_out = (float*)malloc(bytes);
17

18
    // Initialize input array
19
    for (int i = 0; i < size; i++) {
20
        h_in[i] = static_cast<float>(i);
21
    }
22

23
    // Allocate device memory
24
    float *d_in, *d_out;
25
    cudaMalloc(&d_in, bytes);
26
    cudaMalloc(&d_out, bytes);
27

28
    // Copy data to GPU
29
    cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice);
30

31
    // Launch kernel
32
    int blockSize = 256;
33
    int gridSize = (size + blockSize - 1) / blockSize;
34
    doubleArray<<<gridSize, blockSize>>>(d_in, d_out, size);
35

36
    // Copy result back
37
    cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost);
38

39
    // Cleanup
40
    cudaFree(d_in);
41
    cudaFree(d_out);
42
    free(h_in);
43
    free(h_out);
44

45
    return 0;
46
}

Optimizing Memory Usage#

Suppose a more complex kernel benefits from using shared memory. Here’s a toy example of using shared memory to reduce global memory reads when performing a tile-based operation:

1
__global__ void tileKernel(float* d_in, float* d_out, int width) {
2
    // Shared memory tile
3
    __shared__ float tile[256];
4

5
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
6
    // Load data into shared memory (coalesced if consecutive threads read consecutive data)
7
    tile[threadIdx.x] = d_in[globalIdx];
8

9
    // Synchronize to ensure shared memory is fully loaded
10
    __syncthreads();
11

12
    // Perform some operation using tile
13
    float val = tile[threadIdx.x] * 2.0f;
14
    // Store back
15
    d_out[globalIdx] = val;
16
}

By loading the tile into shared memory once, threads can reuse the data multiple times before writing results back, reducing global memory access.

Scheduling Tips and Tricks#

1. Occupancy: Launch enough threads to keep SMs fully utilized. However, remember that extremely large block sizes can reduce the number of concurrent blocks due to register/shared memory limits.
2. Coalescing: Align memory accesses whenever possible to ensure memory transactions are clustered.
3. Minimize Divergence: Warps executing branches cause multiple passes of execution. Structure code to reduce if-else divergence.
4. Overlapping Computation and Memory: In advanced scenarios, you can use CUDA streams to overlap data transfers and kernel executions.

Dynamic Parallelism#

Dynamic parallelism allows a kernel to launch additional kernels from within itself. This is especially useful in scenarios where the amount of parallel work cannot be determined until some partial processing is done (e.g., graph traversal). While this feature can add overhead, it simplifies the programming model for irregular or nested tasks.

GPU Over-subscription#

In large-scale data centers, you may want to run more threads than the GPU can physically handle, with the understanding that the scheduler will context-switch among them. Techniques like NVIDIA’s MPS or AMD Multi-Queue can help multiple processes or multiple kernels to share a single GPU effectively. Memory constraints become critical here, as each process or kernel requires memory resources.

Zero-Copy and Unified Memory Architectures#

• Zero-Copy enables the GPU to access pinned (page-locked) host memory directly, avoiding explicit copying. While it can reduce overhead, it’s typically slower than VRAM for large random accesses, since the data still travels via PCIe or NVLink.
• Unified Memory (or Managed Memory) automates data migration between CPU and GPU memory spaces. This can simplify development but can introduce performance overhead if not used carefully.

Multi-GPU Strategies#

For massive datasets and computations:

• Multi-GPU training or simulation splits the workload among multiple GPUs.
• Techniques like NCCL (NVIDIA Collective Communications Library) or MPI-based methods can handle data exchange between GPUs.
• Schedulers coordinate tasks among GPUs, aiming to maximize parallel usage of each card.