1. What Is GPU Parallelism?
The main idea behind GPU programming is:
Instead of one processor doing one operation at a time, thousands of threads can work on different pieces of data concurrently.
Suppose we have:
1A = [10, 20, 30, 40, 50, 60, 70, 80]
and want:
1B[i] = A[i] × 2
A CPU could process:
1A[0] → ×2 2A[1] → ×2 3A[2] → ×2 4A[3] → ×2 5...
A GPU can conceptually assign:
1Thread 0 → A[0] 2Thread 1 → A[1] 3Thread 2 → A[2] 4Thread 3 → A[3] 5...
So:
1Data 2 ↓ 3Many independent operations 4 ↓ 5Many GPU threads
This is data parallelism.
2. The CUDA Thread Hierarchy
CUDA organizes threads into three major levels:
1Grid 2 ↓ 3Blocks 4 ↓ 5Threads
The hierarchy is:
1GPU Kernel 2 │ 3 ↓ 4 Grid 5 │ 6 ├── Block 0 7 │ ├── Thread 0 8 │ ├── Thread 1 9 │ ├── Thread 2 10 │ └── ... 11 │ 12 ├── Block 1 13 │ ├── Thread 0 14 │ ├── Thread 1 15 │ └── ... 16 │ 17 └── Block 2 18 ├── Thread 0 19 ├── Thread 1 20 └── ...
Remember:
1Grid 2 └── Blocks 3 └── Threads
3. What Is a Thread?
A thread is a logical execution instance of your CUDA kernel.
Suppose:
1__global__ void hello() 2{ 3 printf("Hello\n"); 4}
and:
1hello<<<1, 4>>>();
You launched:
11 block 2× 34 threads
Therefore the kernel executes for:
14 threads
Conceptually:
1Thread 0 → hello() 2Thread 1 → hello() 3Thread 2 → hello() 4Thread 3 → hello()
Every thread executes the same kernel code, but each thread can work on different data.
4. What Is a Block?
A block is a group of threads.
For example:
1hello<<<3, 4>>>();
means:
13 blocks 2× 34 threads/block
Conceptually:
1Grid 2│ 3├── Block 0 4│ ├── Thread 0 5│ ├── Thread 1 6│ ├── Thread 2 7│ └── Thread 3 8│ 9├── Block 1 10│ ├── Thread 0 11│ ├── Thread 1 12│ ├── Thread 2 13│ └── Thread 3 14│ 15└── Block 2 16 ├── Thread 0 17 ├── Thread 1 18 ├── Thread 2 19 └── Thread 3
Total:
13 × 4 = 12 threads
5. What Is a Grid?
The grid is the complete collection of blocks launched for one kernel.
If you write:
1kernel<<<100, 256>>>();
then:
1Grid 2 ├── Block 0 3 ├── Block 1 4 ├── Block 2 5 ├── ... 6 └── Block 99
Each block contains:
1256 threads
Therefore the kernel has:
1100 × 256 2= 25,600 threads
6. The Most Important Formula
To give each thread a unique position in a 1D array:
1int i = 2 blockIdx.x * blockDim.x 3 + threadIdx.x;
This formula connects:
1Grid 2 ↓ 3Block 4 ↓ 5Thread 6 ↓ 7Global data index
Example:
1blockIdx.x = 2 2blockDim.x = 256 3threadIdx.x = 10
Then:
1i = 2 × 256 + 10 2 = 522
So this thread processes:
1data[522]
7. Why Do We Need Blocks?
You might ask:
Why not just launch millions of independent threads?
CUDA groups threads into blocks because threads in the same block can cooperate.
Threads within a block can use:
1Shared memory 2Synchronization 3Cooperative computation
For example:
1Block 2│ 3├── Thread 0 ─┐ 4├── Thread 1 ─┤ 5├── Thread 2 ─┼──→ Shared Memory 6├── Thread 3 ─┤ 7└── ... ─┘
This becomes very important for:
- matrix multiplication
- reductions
- convolution
- attention
- FlashAttention
8. Blocks and SMs
Now connect this to the GPU architecture from Phase 1.
Remember:
1GPU 2 ├── SM 0 3 ├── SM 1 4 ├── SM 2 5 └── ...
When you launch a kernel:
1Grid 2│ 3├── Block 0 ──→ scheduled on an SM 4├── Block 1 ──→ scheduled on an SM 5├── Block 2 ──→ scheduled on an SM 6└── ...
The GPU scheduler decides where blocks execute.
You should not assume:
1Block 0 → SM 0 2Block 1 → SM 1
The actual mapping is handled by the hardware/runtime.
9. Blocks Are Independent
One of the most important CUDA concepts:
Blocks should generally be designed to execute independently.
For example:
1Block 0 2 ↓ 3data 0–255 4 5Block 1 6 ↓ 7data 256–511 8 9Block 2 10 ↓ 11data 512–767
Each block can execute independently.
This allows the GPU to schedule blocks across available SMs.
10. Thousands of Threads
Suppose:
1kernel<<<1024, 256>>>();
Total threads:
11024 × 256 2= 262,144 threads
Conceptually:
1Grid 2│ 3├── Block 0 4│ └── 256 threads 5│ 6├── Block 1 7│ └── 256 threads 8│ 9├── Block 2 10│ └── 256 threads 11│ 12... 13│ 14└── Block 1023 15 └── 256 threads
You have launched:
262,144 logical threads
But your GPU doesn't necessarily execute all 262,144 simultaneously.
This distinction is extremely important.
11. Logical Threads vs Physical Execution
When you launch:
1kernel<<<1024, 256>>>();
you create:
1262,144 logical threads
But the GPU has a finite number of:
1SMs 2execution resources 3registers 4shared memory
Therefore the GPU schedules the work.
Conceptually:
1262,144 logical threads 2 ↓ 3 Blocks 4 ↓ 5 Scheduler 6 ↓ 7 SMs 8 ↓ 9 Warps 10 ↓ 11 Execution
The GPU processes the workload in waves as resources become available.
12. Warps
CUDA threads are executed in groups called warps.
On NVIDIA GPUs:
11 warp = 32 threads
Suppose:
1kernel<<<1, 256>>>();
Then:
1256 / 32 2= 8 warps
Conceptually:
1Block 2│ 3├── Warp 0 → Threads 0–31 4├── Warp 1 → Threads 32–63 5├── Warp 2 → Threads 64–95 6├── Warp 3 → Threads 96–127 7├── Warp 4 → Threads 128–159 8├── Warp 5 → Threads 160–191 9├── Warp 6 → Threads 192–223 10└── Warp 7 → Threads 224–255
This is the foundation of NVIDIA's SIMT execution model.
13. SIMT
SIMT means:
Single Instruction, Multiple Threads
Suppose:
1C[i] = A[i] + B[i];
A warp can conceptually execute:
1Thread 0 → A[0] + B[0] 2Thread 1 → A[1] + B[1] 3Thread 2 → A[2] + B[2] 4... 5Thread 31 → A[31] + B[31]
Same instruction:
1+
Different data:
1A[0], A[1], A[2]...
This is the basic idea behind GPU parallelism.
14. First Working CUDA Program
Let's create a real program that demonstrates:
1Thread 2 ↓ 3Block 4 ↓ 5Grid
thread_block_grid.cu
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4__global__ void show_thread_info() 5{ 6 int global_id = 7 blockIdx.x * blockDim.x 8 + threadIdx.x; 9 10 printf( 11 "Global ID = %d | " 12 "Block = %d | " 13 "Thread = %d\n", 14 global_id, 15 blockIdx.x, 16 threadIdx.x 17 ); 18} 19 20int main() 21{ 22 int blocks = 3; 23 int threads = 4; 24 25 show_thread_info<<<blocks, threads>>>(); 26 27 cudaError_t err = cudaGetLastError(); 28 29 if (err != cudaSuccess) 30 { 31 printf( 32 "Kernel launch error: %s\n", 33 cudaGetErrorString(err) 34 ); 35 36 return 1; 37 } 38 39 err = cudaDeviceSynchronize(); 40 41 if (err != cudaSuccess) 42 { 43 printf( 44 "Kernel execution error: %s\n", 45 cudaGetErrorString(err) 46 ); 47 48 return 1; 49 } 50 51 return 0; 52}
Compile:
1nvcc thread_block_grid.cu -o thread_block_grid
Run:
1./thread_block_grid
You should see output containing values conceptually similar to:
1Global ID = 0 | Block = 0 | Thread = 0 2Global ID = 1 | Block = 0 | Thread = 1 3Global ID = 2 | Block = 0 | Thread = 2 4Global ID = 3 | Block = 0 | Thread = 3 5 6Global ID = 4 | Block = 1 | Thread = 0 7Global ID = 5 | Block = 1 | Thread = 1 8Global ID = 6 | Block = 1 | Thread = 2 9Global ID = 7 | Block = 1 | Thread = 3 10 11Global ID = 8 | Block = 2 | Thread = 0 12Global ID = 9 | Block = 2 | Thread = 1 13Global ID = 10 | Block = 2 | Thread = 2 14Global ID = 11 | Block = 2 | Thread = 3
Important: CUDA thread execution is not guaranteed to print in this neat order. printf output ordering is not something you should rely on.
15. Real Example — Parallel Array Processing
Now let's do something useful.
Suppose:
1A = [1, 2, 3, 4, 5, 6, 7, 8]
We want:
1B = A × 10
Mathematically:
1B[i] = A[i] × 10
CUDA:
1__global__ void multiply_by_10( 2 const float* A, 3 float* B, 4 int N 5) 6{ 7 int i = 8 blockIdx.x * blockDim.x 9 + threadIdx.x; 10 11 if (i < N) 12 { 13 B[i] = A[i] * 10.0f; 14 } 15}
Each thread processes one element:
1Thread 0 → A[0] 2Thread 1 → A[1] 3Thread 2 → A[2] 4...
16. Complete Working Vector Example
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4__global__ void multiply_by_10( 5 const float* A, 6 float* B, 7 int N 8) 9{ 10 int i = 11 blockIdx.x * blockDim.x 12 + threadIdx.x; 13 14 if (i < N) 15 { 16 B[i] = A[i] * 10.0f; 17 } 18} 19 20int main() 21{ 22 const int N = 16; 23 24 float h_A[N]; 25 float h_B[N]; 26 27 for (int i = 0; i < N; i++) 28 { 29 h_A[i] = static_cast<float>(i + 1); 30 } 31 32 float* d_A; 33 float* d_B; 34 35 size_t bytes = N * sizeof(float); 36 37 cudaMalloc(&d_A, bytes); 38 cudaMalloc(&d_B, bytes); 39 40 cudaMemcpy( 41 d_A, 42 h_A, 43 bytes, 44 cudaMemcpyHostToDevice 45 ); 46 47 int threads = 256; 48 49 int blocks = 50 (N + threads - 1) / threads; 51 52 multiply_by_10<<<blocks, threads>>>( 53 d_A, 54 d_B, 55 N 56 ); 57 58 cudaError_t err = cudaGetLastError(); 59 60 if (err != cudaSuccess) 61 { 62 printf( 63 "Launch error: %s\n", 64 cudaGetErrorString(err) 65 ); 66 67 return 1; 68 } 69 70 err = cudaDeviceSynchronize(); 71 72 if (err != cudaSuccess) 73 { 74 printf( 75 "Execution error: %s\n", 76 cudaGetErrorString(err) 77 ); 78 79 return 1; 80 } 81 82 cudaMemcpy( 83 h_B, 84 d_B, 85 bytes, 86 cudaMemcpyDeviceToHost 87 ); 88 89 printf("Results:\n"); 90 91 for (int i = 0; i < N; i++) 92 { 93 printf( 94 "A[%d] = %.1f -> B[%d] = %.1f\n", 95 i, 96 h_A[i], 97 i, 98 h_B[i] 99 ); 100 } 101 102 cudaFree(d_A); 103 cudaFree(d_B); 104 105 return 0; 106}
Compile:
1nvcc multiply.cu -o multiply
Run:
1./multiply
Expected result:
1A[0] = 1.0 -> B[0] = 10.0 2A[1] = 2.0 -> B[1] = 20.0 3A[2] = 3.0 -> B[2] = 30.0 4... 5A[15] = 16.0 -> B[15] = 160.0
17. What Actually Happened?
You wrote:
1multiply_by_10<<<blocks, threads>>>( 2 d_A, 3 d_B, 4 N 5);
Suppose:
1N = 16 2threads = 256
Then:
1blocks = 1
So:
1Grid 2└── Block 0 3 ├── Thread 0 → B[0] 4 ├── Thread 1 → B[1] 5 ├── Thread 2 → B[2] 6 ├── ... 7 ├── Thread 15 → B[15] 8 └── Threads 16–255 → do nothing
Because:
1if (i < N)
prevents invalid accesses.
18. Scaling to 1 Million Elements
Now change:
1const int N = 1000000;
You don't need to rewrite the kernel.
The same kernel can process one million elements:
11,000,000 elements 2 ↓ 3many blocks 4 ↓ 5many threads 6 ↓ 7GPU
This is the power of CUDA's execution model.
The kernel is written once:
1B[i] = A[i] * 10;
and CUDA maps the work across the GPU.
19. Why This Is Powerful for AI
Consider an AI tensor with millions of values.
For example:
1Activation 2Shape: 3[4096, 4096]
Number of elements:
14096 × 4096 2= 16,777,216
A simple elementwise operation such as:
1Y = SiLU(X)
can conceptually assign one thread per element:
1Thread 0 → X[0] 2Thread 1 → X[1] 3Thread 2 → X[2] 4...
This is why GPU parallelism is fundamental to deep learning.
20. Example: ReLU Kernel
Mathematical definition:
1ReLU(x) = max(0, x)
CUDA:
1__global__ void relu( 2 const float* input, 3 float* output, 4 int N 5) 6{ 7 int i = 8 blockIdx.x * blockDim.x 9 + threadIdx.x; 10 11 if (i < N) 12 { 13 float x = input[i]; 14 15 output[i] = 16 (x > 0.0f) ? x : 0.0f; 17 } 18}
The mapping is:
1Thread 2 ↓ 3Global index 4 ↓ 5Tensor element 6 ↓ 7ReLU 8 ↓ 9Output tensor
21. Example: SiLU Kernel
The SiLU/Swish activation is:
1SiLU(x) = x × sigmoid(x)
where:
1sigmoid(x) = 1 / (1 + e^-x)
CUDA:
1__global__ void silu( 2 const float* input, 3 float* output, 4 int N 5) 6{ 7 int i = 8 blockIdx.x * blockDim.x 9 + threadIdx.x; 10 11 if (i < N) 12 { 13 float x = input[i]; 14 15 float sigmoid = 16 1.0f / (1.0f + expf(-x)); 17 18 output[i] = 19 x * sigmoid; 20 } 21}
This is a very simple example of an AI-oriented GPU kernel.
22. Thread-Level Parallelism vs Block-Level Parallelism
There are two useful ways to think about the workload.
Thread-level
1Thread 0 → element 0 2Thread 1 → element 1 3Thread 2 → element 2
Block-level
1Block 0 → elements 0–255 2Block 1 → elements 256–511 3Block 2 → elements 512–767
Together:
1Grid 2│ 3├── Block 0 → Threads → Data 4├── Block 1 → Threads → Data 5├── Block 2 → Threads → Data 6└── ...
23. Why Millions of Threads Don't Mean Millions of Cores
This is a very important concept.
Suppose you launch:
110 million threads
Your GPU does not need:
110 million physical cores
Instead:
1Millions of logical threads 2 ↓ 3Blocks 4 ↓ 5Scheduler 6 ↓ 7Available SM resources 8 ↓ 9Warps 10 ↓ 11Execution
The GPU keeps work available and schedules it as hardware resources permit.
This is one of the fundamental differences between the CUDA programming model and the physical GPU hardware.
24. Thread Independence
The easiest CUDA kernels are embarrassingly parallel.
For example:
1B[i] = A[i] × 2
Thread 0 doesn't need Thread 1.
1Thread 0 → A[0] 2Thread 1 → A[1] 3Thread 2 → A[2]
They can work independently.
This is ideal for GPU execution.
25. When Threads Need Cooperation
Some algorithms require threads to share information.
For example:
1Sum: 21 + 2 + 3 + 4 + ... + N
You cannot simply have:
1Thread 0 → result 2Thread 1 → result 3Thread 2 → result
without coordination.
You need:
1Threads 2 ↓ 3Partial results 4 ↓ 5Synchronization 6 ↓ 7Reduction 8 ↓ 9Final result
This leads to the next important CUDA topics:
1Shared Memory 2 ↓ 3__syncthreads() 4 ↓ 5Reduction 6 ↓ 7Warp-level primitives
26. Thread/Block/Grid vs GPU Architecture
Connect everything you've learned:
1CUDA Software Model 2 │ 3 ↓ 4 Grid 5 │ 6 ↓ 7 Blocks 8 │ 9 ↓ 10 Threads 11 │ 12 ↓ 13 Warps 14 │ 15 ↓ 16GPU Hardware 17 │ 18 ↓ 19 SMs 20 │ 21 ├── Registers 22 ├── Shared Memory 23 ├── CUDA Cores 24 └── Tensor Cores
This is the mental model you need for kernel optimization.
27. Important Relationship
Don't think:
1Thread = CUDA Core
That's incorrect.
Instead think:
1Logical CUDA threads 2 ↓ 3Grouped into warps 4 ↓ 5Scheduled on SMs 6 ↓ 7Execute using available hardware resources
This distinction becomes extremely important when you study:
- occupancy
- register pressure
- warp scheduling
- latency hiding
- Tensor Cores
- kernel optimization
28. Practical Exercise
Try changing:
1int blocks = 3; 2int threads = 4;
to:
1int blocks = 4; 2int threads = 8;
Now calculate:
14 × 8 2= 32 threads
Then try:
1int blocks = 16; 2int threads = 256;
Calculate:
116 × 256 2= 4096 threads
Then:
1int blocks = 1024; 2int threads = 256;
Calculate:
11024 × 256 2= 262,144 threads
This exercise helps you develop an intuition for CUDA's execution model.
29. Final Mental Model
You should now be able to visualize:
1 GPU 2 │ 3 ↓ 4 Kernel 5 │ 6 ↓ 7 Grid 8 │ 9 ┌─────────┼─────────┐ 10 ↓ ↓ ↓ 11 Block 0 Block 1 Block 2 12 │ │ │ 13 Threads Threads Threads 14 │ │ │ 15 Warps Warps Warps 16 │ │ │ 17 └─────────┼─────────┘ 18 ↓ 19 SMs 20 │ 21 ┌────────────┼────────────┐ 22 ↓ ↓ ↓ 23 Registers Shared Mem Execution
And for a 1D problem:
1int i = 2 blockIdx.x * blockDim.x 3 + threadIdx.x;
means:
1Which block am I in? 2 ↓ 3How large is each block? 4 ↓ 5Which thread am I? 6 ↓ 7 i 8 ↓ 9Which data element should I process?
Your Phase 2 milestone
You should be able to write this without looking at notes:
1__global__ void kernel( 2 float* data, 3 int N 4) 5{ 6 int i = 7 blockIdx.x * blockDim.x 8 + threadIdx.x; 9 10 if (i < N) 11 { 12 data[i] *= 2.0f; 13 } 14}
and launch it with:
1int threads = 256; 2 3int blocks = 4 (N + threads - 1) / threads; 5 6kernel<<<blocks, threads>>>(data, N);
If you understand why every line exists, you have a solid foundation for the next stage: warps, SIMT, warp divergence, synchronization, shared memory, and memory coalescing.