GPUS and How to Program Them by Manya Bansal

CUDA, designed as an extension to C++, preserves its familiar abstractions. However, unlike CPU programming --- where compilers and runtime systems abstract away most hardware concerns --- writing CUDA code requires developers to manually map computations onto the GPU’s parallel execution and memory hierarchy, while respecting the fundamental constraints of the hardware. In this talk, I’ll discuss where this model breaks down and why an alternative low-level language is needed for productive, compositional GPU programming.

1. A ScyllaDB Community
GPUS, and How to Program
Them
Manya Bansal
Ph.D. Student

2. A ScyllaDB Community
GPUS, and how to program
them
Manya Bansal
Ph.D. Student
not
^

3. The Era of Compute

4. Era of GPUs
Source: Dally, B. (2023, August 29). Hardware for Deep Learning. Hot Chips 35 Symposium.
K20X M40 V100 A100 H100

5. Explosion in
the amount
of compute
Source: Dally, B. (2023, August 29). Hardware for Deep Learning. Hot Chips 35
Symposium. Retrieved from Hot Chips website:

6. H100

7. GPU’s Hardware

8. CPU Execution Model

9. CPU Execution Model
ADD R0, R0, #1
CMP R0, R1
BNE loop
SUB R2, R1, #2
MOV R7, #1
MOV R0, #0
SWI 0

10. CPU Execution Model
ADD R0, R0, #1
CMP R0, R1
BNE loop
SUB R2, R1, #2
MOV R7, #1
MOV R0, #0
SWI 0
Compiler
Operating System
Tower of Abstractions

11. GPU Execution Model
…
Warp

12. …
Warp
Streaming Multiprocessor
…
Warp
…
Warp
…
Warp

13. …
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
GPU

14. GPU Execution Model
…
SM
Instruction Cache
Warp
Scheduler
Warp
Scheduler
Register File
LD/ST LD/STLD/ST
Tensor CoresFP32

15. Deep Memory Hiearchy
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
GPU
Global Memory

16. Deep Memory Hiearchy
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
GPU
Global Memory
…
Warp
Threads in a Block
…
Warp
…
Warp
…
Warp
Shared Memory

17. Deep Memory Hiearchy
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
…
Warp
Simultaneous Multiprocessor
…
Warp
…
Warp
…
Warp
GPU
Global Memory
…
Warp
Threads in a Block
…
Warp
…
Warp
…
Warp
Shared Memory
Registers

18. SM
L2
Cache
HBM
Memory
Controller

19. CUDA: How we program GPUs today

20. “Basically”, C++

21. A Simple Program
a
1
+
b

22. __device__ void add_1_scalar(const float* a,
float* b, int N){
for (int i 0; i < N; i++)
b[i] = a[i] + 1;
}

__global__ void add_1_scalar_kernel(
const float *d_a,
float *d_b, int N){
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N){
int elements_to_process = min(4, N - tid);
add_1_scalar(&d_a[tid], &d_b[tid],
elements_to_process);
}}


… ……

23. __device__ void add_1_thread(const float* a,
float* b, int N){
if (threadIdx.x < N)
b[threadIdx.x] = a[threadIdx.x] + 1;}

__global__ void add_1_thread_kernel(
const float* d_a,
float* d_b, int total_N){
int block_offset = blockIdx.x * blockDim.x;
int elements_in_block = min(blockDim.x,
total_N - block_offset);
if (block_offset < total_N) {
add_1_thread(&d_a[block_offset],
& d_b[block_offset],
elements_in_block);
}}

… ……

24. __device__ void add_1_shared(const float* a,
float* b, int N){
extern __shared__ float smem[];
if (threadIdx.x < N)
smem[(threadIdx.x + 1) % N] = a[(threadIdx.x + 1) % N];
__syncthreads();
if (threadIdx.x < N)
b[threadIdx.x] = smem[threadIdx.x] + 1; }

__global__ void add_1_thread_kernel(
const float* d_a,
float* d_b, int total_N){
int block_offset = blockIdx.x * blockDim.x;
int elements_in_block = min(blockDim.x,
total_N - block_offset);
if (block_offset < total_N) {
add_1_thread(&d_a[block_offset],&d_b[block_offset],
elements_in_block); }}

… ……

25. “Basically”, C++
__device__ void add_1_scalar(const float* a,
float* b, int N)
__device__ void add_1_thread(const float* a,
float* b, int N);
__device__ void add_1_shared(const float* a,
float* b, int N);

26. Lack of composition has fragmented the ecosystem

27. Hardware is missing!

28. Programming GPUs uses finite
resources, like memory and compute.

29. Model the compute
hierarchy explicitly
__device__ void add_1(block : 1 const float *a,
block : 1 float *b, int N)
{
at(block) blk;
invoke(blk)
{
if (__threadId < N)
b[__threadId] = a[__threadId] + 1;
}
}

30. Structured
Decomposition of
Parallel work

__device__ void add_1_thread(
block:1 const float* a,
block:1 float* b, int N){
at(block:2) b2;
invoke(b2) {
at(block:1) b1;
invoke(b1) {
if (__threadId < N)
b[__threadId] = a[__threadId] + 1;
}
}
}

31. Shared memory is
also special syntax
at (block:1) b;
invoke(b) {
SHARED float shared_array[4][8];
}

32. Names and
invocation lead to
a natural pattern
for concurrent and
parallel execution
at (thread:2) t1;
at (thread:2) t2;

invoke(t1) { ... }

invoke(t2) { ... }

invoke(t1) { ... }

33. Pointers are also
tied to compute
resources
__device__ void add_1_thread(
const block : 1 float *a,
block : 1 float *b, int N)
{
at(block : 1) b;
invoke(b)
{
at(thread : 1) t;
invoke(t)
{
// Error
b[0] = 0;
}}}

34. Memory can be
lowered to access
different parts of
the compute
hierarchy
__device__ void add_1_thread(
const block : 1 float *a,
block : 1 float *b, int N)
{
at(block : 1) b;
invoke(b)
{
at(thread : 1) t;
thread:
1 float *bt<> b[__thread_id];
invoke(t)
{
// bt[i] = b[__thread_id + i]
bt[0] = ...;
}}}

35. Composition
becomes a problem
of deciding whether
the required
resources can be
provided
__device__ void add_1_thread(
const block:1 float* a,
block:1 float* b, int N)
@ {block>=1, thread>=1, SMEM: 1024B};

36. Low-level, safe
substrate makes it
easier to build
productive tools

37. Safe and compositional CUDA
through scoped resources
[email protected]://manya-bansal.github.io/