1. Why GPU Memory Matters
A CUDA kernel can have perfect mathematical logic and still be slow.
One major reason:
The kernel spends too much time waiting for data.
Consider:
1GPU computation 2 ↓ 3Need data 4 ↓ 5Memory access 6 ↓ 7Data arrives 8 ↓ 9Compute
If memory is slow relative to computation:
1Compute → waiting → waiting → waiting
The GPU's arithmetic hardware may be underutilized.
Therefore, high-performance CUDA programming requires understanding:
1Where is my data? 2 ↓ 3How far away is it? 4 ↓ 5How fast can I access it? 6 ↓ 7How do neighboring threads access it?
2. CUDA Memory Hierarchy
A simplified NVIDIA GPU memory hierarchy looks like:
1 GPU 2 │ 3 ↓ 4 SM 5 │ 6 ┌────────────┼────────────┐ 7 ↓ ↓ ↓ 8 Registers Shared Memory L1 Cache 9 │ │ 10 └────────────┼────────────┘ 11 ↓ 12 L2 Cache 13 ↓ 14 Global Memory
There are also other memory spaces, such as:
1Constant Memory 2Texture Memory 3Local Memory
For your kernel-programming path, initially focus on:
11. Registers 22. Shared Memory 33. L1 Cache 44. L2 Cache 55. Global Memory
3. The Basic Rule
Generally:
1Closer to computation 2 ↓ 3Lower access latency 4 ↓ 5Higher potential performance
A simplified conceptual ordering is:
1Registers 2 ↓ 3Shared Memory / L1 4 ↓ 5L2 6 ↓ 7Global Memory
But don't memorize exact latency numbers.
They depend on:
1GPU architecture 2access pattern 3cache state 4contention 5compiler behavior
The important concept is the hierarchy.
4. Registers
Registers are the fastest storage available to a CUDA thread.
Think:
1Thread 2 │ 3 └── Registers
Each thread has its own register state.
For example:
1__global__ void example() 2{ 3 int a = 10; 4 int b = 20; 5 6 int c = a + b; 7}
The compiler may place these values in registers.
Conceptually:
1Thread 2 ├── Register → a 3 ├── Register → b 4 └── Register → c
5. Register Scope
Registers are private to each thread.
For example:
1Thread 0 2 ├── R0 3 ├── R1 4 └── R2 5 6Thread 1 7 ├── R0 8 ├── R1 9 └── R2
Thread 0 cannot directly access Thread 1's registers.
This is useful for temporary values.
6. Register Pressure
Registers are fast, but they are limited.
Suppose a kernel requires many temporary variables:
1float a1; 2float a2; 3float a3; 4float a4; 5float a5; 6float a6; 7...
The compiler may need many registers per thread.
If register demand becomes too high, the compiler can spill some values to local memory.
Conceptually:
1Too many registers 2 ↓ 3Register pressure 4 ↓ 5Spilling 6 ↓ 7Local-memory accesses 8 ↓ 9Potential performance loss
This becomes very important when optimizing large AI kernels.
7. Shared Memory
Shared memory is memory located on/near each SM and shared among threads in a block.
Conceptually:
1SM 2│ 3├── Warp 0 4├── Warp 1 5├── Warp 2 6│ 7└── Shared Memory 8 ↑ 9 │ 10 Threads in block
Unlike registers:
1Register 2→ private to thread
Shared memory:
1Shared memory 2→ accessible by threads in the block
8. Declaring Shared Memory
Static shared memory:
1__shared__ float data[256];
Example:
1__global__ void example() 2{ 3 __shared__ float data[256]; 4 5 int tid = threadIdx.x; 6 7 data[tid] = tid; 8}
Each block gets its own instance.
If you launch:
1Block 0 2Block 1 3Block 2
conceptually:
1Block 0 → Shared Memory A 2Block 1 → Shared Memory B 3Block 2 → Shared Memory C
Blocks do not share the same block-local shared-memory array.
9. Shared Memory Is Block-Scoped
This is extremely important.
Suppose:
1__shared__ float data[256];
Threads in Block 0 can cooperate through:
1Block 0 2 ├── Thread 0 ──┐ 3 ├── Thread 1 ──┤ 4 ├── Thread 2 ──┼──→ data[] 5 └── ... ┘
But Block 1 has a separate:
1data[]
So:
1Block 0 → shared memory instance 0 2 3Block 1 → shared memory instance 1
10. Shared Memory + Synchronization
This connects directly to your previous lesson.
Suppose:
1data[tid] = input[tid];
Then:
1__syncthreads();
Then:
1use data[]
The pattern is:
1Global Memory 2 ↓ 3Load 4 ↓ 5Shared Memory 6 ↓ 7__syncthreads() 8 ↓ 9Reuse data 10 ↓ 11Compute
This is one of the most important CUDA optimization patterns.
11. Why Shared Memory?
Suppose 32 threads repeatedly need the same data.
Without effective reuse:
1Thread 0 → Global Memory 2Thread 1 → Global Memory 3Thread 2 → Global Memory 4...
With shared-memory tiling:
1Global Memory 2 ↓ 3Load once 4 ↓ 5Shared Memory 6 ↓ 7Multiple threads reuse
The idea is:
Move data closer to computation and reuse it.
This is called data reuse.
12. Simple Shared Memory Example
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4__global__ void shared_example( 5 const float* input, 6 float* output 7) 8{ 9 __shared__ float tile[256]; 10 11 int tid = threadIdx.x; 12 13 tile[tid] = input[tid]; 14 15 __syncthreads(); 16 17 output[tid] = 18 tile[tid] * 2.0f; 19} 20 21int main() 22{ 23 const int N = 256; 24 25 float h_input[N]; 26 float h_output[N]; 27 28 for (int i = 0; i < N; i++) 29 { 30 h_input[i] = static_cast<float>(i); 31 } 32 33 float* d_input; 34 float* d_output; 35 36 size_t bytes = 37 N * sizeof(float); 38 39 cudaMalloc( 40 &d_input, 41 bytes 42 ); 43 44 cudaMalloc( 45 &d_output, 46 bytes 47 ); 48 49 cudaMemcpy( 50 d_input, 51 h_input, 52 bytes, 53 cudaMemcpyHostToDevice 54 ); 55 56 shared_example<<<1, 256>>>( 57 d_input, 58 d_output 59 ); 60 61 cudaError_t err = 62 cudaGetLastError(); 63 64 if (err != cudaSuccess) 65 { 66 printf( 67 "Launch error: %s\n", 68 cudaGetErrorString(err) 69 ); 70 71 return 1; 72 } 73 74 cudaDeviceSynchronize(); 75 76 cudaMemcpy( 77 h_output, 78 d_output, 79 bytes, 80 cudaMemcpyDeviceToHost 81 ); 82 83 for (int i = 0; i < 10; i++) 84 { 85 printf( 86 "%.1f ", 87 h_output[i] 88 ); 89 } 90 91 printf("\n"); 92 93 cudaFree(d_input); 94 cudaFree(d_output); 95 96 return 0; 97}
Expected beginning:
10.0 2.0 4.0 6.0 8.0 ...
13. Important: This Example Isn't Automatically Faster
This is a very important optimization lesson.
You might think:
1Global → Shared → Compute
must always be faster.
Not necessarily.
In our example:
1tile[tid] = input[tid]; 2 3__syncthreads(); 4 5output[tid] = tile[tid] * 2;
Each value is loaded once and used once.
The shared-memory copy may actually add overhead.
Shared memory becomes especially useful when:
1Data is reused
multiple times.
14. Data Reuse
Imagine:
1A value 2 ↓ 3used 10 times
Global-memory approach:
1Load 2Compute 3Load 4Compute 5Load 6Compute 7...
Shared-memory approach:
1Global Load 2 ↓ 3Shared Memory 4 ↓ 5Reuse 6Reuse 7Reuse 8Reuse 9...
Now the investment of loading into shared memory can make more sense.
15. Global Memory
Global memory is the large memory accessible by GPU threads.
It is typically backed by the GPU's VRAM.
For example:
1float* d_data; 2cudaMalloc(&d_data, bytes);
The pointer:
1d_data
points to device memory.
Conceptually:
1GPU VRAM 2└── Global Memory 3 ├── Tensor A 4 ├── Tensor B 5 ├── Tensor C 6 └── ...
Global memory has high capacity compared with registers/shared memory, but access latency is much higher than registers.
16. L1 Cache
L1 cache is a small, fast cache associated with the SM.
Conceptually:
1SM 2│ 3├── Registers 4├── Shared Memory 5└── L1 Cache
The exact relationship between L1 and shared memory depends on GPU architecture and configuration.
For programming, the important idea is:
Frequently accessed global-memory data may be served from cache rather than going all the way to VRAM.
17. L2 Cache
L2 is a larger cache shared across the GPU.
Conceptually:
1SM 0 ──┐ 2SM 1 ──┤ 3SM 2 ──┼──→ L2 Cache 4SM 3 ──┘ 5 ↓ 6 Global Memory
This means multiple SMs can benefit from data that is resident in L2.
18. Simplified Memory Hierarchy
For learning, remember:
1 GPU 2 │ 3 SM 4 │ 5 ┌──────┼──────┐ 6 ↓ ↓ ↓ 7 Registers Shared L1 8 Memory 9 │ │ │ 10 └──────┼──────┘ 11 ↓ 12 L2 13 ↓ 14 Global Memory
And think:
1Registers 2 ↓ 3very fast / very limited / thread-private 4 5Shared Memory 6 ↓ 7fast / limited / block-shared 8 9L1 10 ↓ 11small cache 12 13L2 14 ↓ 15larger cache 16 17Global Memory 18 ↓ 19large / high latency
19. The Most Important Optimization Question
When optimizing a kernel, don't ask only:
"How many FLOPs does this kernel perform?"
Also ask:
"How many times am I moving the same data?"
For example:
1Load A 2Load B 3Compute 4Load A again 5Load B again 6Compute
might be improved through reuse:
1Load A → Shared Memory 2Load B → Shared Memory 3 ↓ 4Compute 5Compute 6Compute
This is the foundation of tiling.
20. Matrix Multiplication
This is where shared memory becomes extremely important.
Suppose:
1C = A × B
Naively, many threads may repeatedly read values from global memory.
A tiled implementation instead loads blocks of matrices into shared memory.
Conceptually:
1Global A ──────┐ 2 ↓ 3 Shared A Tile 4 │ 5 ↓ 6 Compute 7 ↑ 8 Shared B Tile 9 ↑ 10Global B ──────┘
21. Tiling
Suppose:
1A = large matrix 2B = large matrix
Instead of processing the entire matrix at once:
1Large Matrix
divide it into tiles:
1┌────┬────┬────┬────┐ 2│ T0 │ T1 │ T2 │ T3 │ 3├────┼────┼────┼────┤ 4│ T4 │ T5 │ T6 │ T7 │ 5├────┼────┼────┼────┤ 6│... │... │... │... │ 7└────┴────┴────┴────┘
Load a tile:
1Global 2 ↓ 3Shared 4 ↓ 5Reuse
22. Simplified Tiled Matrix Multiplication
Here's a learning implementation.
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4#define TILE 16 5 6__global__ void matmul( 7 const float* A, 8 const float* B, 9 float* C, 10 int N 11) 12{ 13 __shared__ float As[TILE][TILE]; 14 __shared__ float Bs[TILE][TILE]; 15 16 int row = 17 blockIdx.y * TILE 18 + threadIdx.y; 19 20 int col = 21 blockIdx.x * TILE 22 + threadIdx.x; 23 24 float sum = 0.0f; 25 26 for (int tile = 0; 27 tile < N; 28 tile += TILE) 29 { 30 int A_col = 31 tile + threadIdx.x; 32 33 int B_row = 34 tile + threadIdx.y; 35 36 if (row < N && A_col < N) 37 { 38 As[threadIdx.y][threadIdx.x] = 39 A[row * N + A_col]; 40 } 41 else 42 { 43 As[threadIdx.y][threadIdx.x] = 44 0.0f; 45 } 46 47 if (B_row < N && col < N) 48 { 49 Bs[threadIdx.y][threadIdx.x] = 50 B[B_row * N + col]; 51 } 52 else 53 { 54 Bs[threadIdx.y][threadIdx.x] = 55 0.0f; 56 } 57 58 __syncthreads(); 59 60 for (int k = 0; 61 k < TILE; 62 k++) 63 { 64 sum += 65 As[threadIdx.y][k] * 66 Bs[k][threadIdx.x]; 67 } 68 69 __syncthreads(); 70 } 71 72 if (row < N && col < N) 73 { 74 C[row * N + col] = sum; 75 } 76}
This kernel demonstrates the core pattern:
1Global Memory 2 ↓ 3Shared Memory 4 ↓ 5__syncthreads() 6 ↓ 7Reuse tile 8 ↓ 9Compute 10 ↓ 11__syncthreads() 12 ↓ 13Next tile
23. Why Are There Two __syncthreads() Calls?
This is very important.
First barrier
1__syncthreads();
means:
Everyone has finished loading the current tiles before anyone uses them.
Then:
1Shared A 2Shared B 3 ↓ 4Compute
Second barrier
1__syncthreads();
means:
Everyone has finished using the current tiles before any thread overwrites them with the next tiles.
So:
1Load 2 ↓ 3SYNC 4 ↓ 5Compute 6 ↓ 7SYNC 8 ↓ 9Load next tile
This pattern appears throughout optimized CUDA programming.
24. Memory Coalescing
Now we reach another major topic.
Suppose 32 threads access:
1data[0] 2data[1] 3data[2] 4... 5data[31]
This is a contiguous access pattern.
Conceptually:
1Thread 0 → data[0] 2Thread 1 → data[1] 3Thread 2 → data[2] 4... 5Thread 31 → data[31]
This is generally favorable for global-memory throughput.
This concept is called:
Memory coalescing
25. Bad Access Pattern
Consider:
1Thread 0 → data[0] 2Thread 1 → data[32] 3Thread 2 → data[64] 4Thread 3 → data[96] 5...
Now threads access memory with a large stride.
Conceptually:
1T0 → 0 2T1 → 32 3T2 → 64 4T3 → 96
This can be much less efficient depending on the access pattern and architecture.
26. Why Coalescing Matters
GPU memory systems move data in transactions.
You want neighboring threads to access neighboring memory locations when possible.
Good:
1Thread → Memory 2 30 → 0 41 → 1 52 → 2 63 → 3 74 → 4 8...
Bad:
1Thread → Memory 2 30 → 0 41 → 1000 52 → 2000 63 → 3000
The exact transaction behavior depends on the GPU architecture and data type, but the general optimization principle is:
Make neighboring threads access nearby data whenever the algorithm allows it.
27. Shared Memory Bank Conflicts
Shared memory has another performance issue:
Bank conflicts
Shared memory is divided into banks.
Conceptually:
1Shared Memory 2│ 3├── Bank 0 4├── Bank 1 5├── Bank 2 6├── ... 7└── Bank N
If threads in a warp access different banks:
1Thread 0 → Bank 0 2Thread 1 → Bank 1 3Thread 2 → Bank 2 4Thread 3 → Bank 3
this is generally favorable.
But if many threads contend for the same bank:
1Thread 0 ─┐ 2Thread 1 ─┤ 3Thread 2 ─┼──→ same bank 4Thread 3 ─┤ 5...
a bank conflict may occur, reducing performance.
The exact behavior has special cases, including broadcast patterns, so don't reduce the concept to "same bank is always bad."
28. Three Major Memory Optimization Concepts
You should now remember:
1Global Memory 2 ↓ 3Memory Coalescing
and:
1Shared Memory 2 ↓ 3Bank Conflicts
and:
1Registers 2 ↓ 3Register Pressure
These three concepts become extremely important in kernel optimization.
29. CUDA Memory Strategy for AI
For a Transformer-style kernel, think:
1Large Model Tensor 2 ↓ 3Global Memory 4 ↓ 5Coalesced Loads 6 ↓ 7Shared Memory / Cache 8 ↓ 9Registers 10 ↓ 11Compute 12 ↓ 13Output
The goal isn't necessarily to manually move everything into shared memory.
Instead:
Use the memory hierarchy intelligently based on reuse, access patterns, and resource limits.
30. Memory vs Compute
A kernel can be limited by:
Compute
1Lots of arithmetic 2 ↓ 3ALU/Tensor Core utilization 4 ↓ 5Compute-bound
or:
Memory
1Lots of data movement 2 ↓ 3Memory system bottleneck 4 ↓ 5Memory-bound
This distinction is fundamental.
For example:
1Simple vector addition 2A[i] + B[i]
does relatively little computation per byte loaded.
It can be memory-bandwidth sensitive.
Whereas:
1Matrix multiplication
can perform enormous amounts of arithmetic while reusing data.
31. Arithmetic Intensity
A useful concept is:
1Arithmetic Intensity 2= 3Number of operations 4-------------------- 5Bytes moved
High arithmetic intensity:
1many computations 2per byte loaded
Low arithmetic intensity:
1few computations 2per byte loaded
This helps reason about whether a kernel may be:
1memory-bound
or:
1compute-bound
32. AI Example
Consider:
1Y = X × W
Matrix multiplication can reuse values of:
1X 2W
across many output computations.
Therefore optimized kernels try to maximize reuse through:
1Registers 2Shared memory 3Caches 4Tensor Cores
This is one reason high-performance GEMM kernels are much more complicated than:
1C[i] = A[i] * B[i];
33. What You Should Memorize
Registers
1Private to thread 2Very fast 3Limited
Shared Memory
1Shared within block 2Fast 3Limited 4Useful for data reuse
L1
1Small cache near SM
L2
1Larger cache shared across GPU
Global Memory
1Large capacity 2High latency compared with on-chip storage
34. The Complete Mental Model
You should now visualize:
1 GPU 2 │ 3 ┌───────┴───────┐ 4 ↓ ↓ 5 SM 0 SM 1 6 │ │ 7 ┌────┼────┐ ┌────┼────┐ 8 ↓ ↓ ↓ ↓ ↓ ↓ 9 Registers Shared L1 Registers Shared L1 10 │ │ 11 └───────┬───────┘ 12 ↓ 13 L2 14 ↓ 15 Global Memory
And the optimization path:
1Global Memory 2 ↓ 3Coalesced access 4 ↓ 5Cache / Shared Memory 6 ↓ 7Register reuse 8 ↓ 9Compute
35. Connection to Your Previous Topics
You now have:
1Threads 2 ↓ 3Blocks 4 ↓ 5Warps 6 ↓ 7SIMT 8 ↓ 9Synchronization 10 ↓ 11Shared Memory
The next major step is:
1Memory Access Patterns 2 ↓ 3Coalesced Memory Access 4 ↓ 5Shared Memory Bank Conflicts 6 ↓ 7Memory Bandwidth 8 ↓ 9Roofline / Arithmetic Intensity
This is where you start moving from "I can write CUDA kernels" toward "I can optimize CUDA kernels."
Practice Challenge
Before moving forward, write a CUDA kernel that performs:
1C[i] = A[i] + B[i]
and then modify it to use:
1__shared__ float tileA[256]; 2__shared__ float tileB[256];
The goal isn't necessarily to make the second version faster. Instead, understand the complete data path:
1CPU 2 ↓ 3Global Memory 4 ↓ 5Shared Memory 6 ↓ 7__syncthreads() 8 ↓ 9Registers 10 ↓ 11Arithmetic 12 ↓ 13Global Memory 14 ↓ 15CPU
Once you understand that path, the next topic—CUDA Memory Coalescing and Access Patterns—is one of the highest-value topics you can learn for real kernel optimization.