1. What Is Memory Coalescing?
Memory coalescing is a fundamental CUDA optimization technique where threads in a warp access memory in a pattern that allows the GPU to service those accesses efficiently.
The simplest mental model is:
1Warp 2 │ 3 ├── Thread 0 → data[0] 4 ├── Thread 1 → data[1] 5 ├── Thread 2 → data[2] 6 ├── Thread 3 → data[3] 7 ├── ... 8 └── Thread 31 → data[31]
The threads access neighboring elements.
This is generally a good global-memory access pattern.
2. Why Does This Matter?
A GPU doesn't conceptually perform:
1Thread 0 → one completely independent memory operation 2Thread 1 → another completely independent memory operation 3Thread 2 → ...
The memory system services accesses through memory transactions.
Therefore, you want the accesses generated by a warp to use memory efficiently.
Think:
1Bad: 2Warp 3 ↓ 4scattered addresses 5 ↓ 6many memory transactions 7 ↓ 8lower effective bandwidth
versus:
1Good: 2Warp 3 ↓ 4nearby addresses 5 ↓ 6efficient transactions 7 ↓ 8higher effective bandwidth
The exact transaction behavior depends on GPU architecture, data type, alignment, cache state, and access pattern.
3. First Example: Perfectly Contiguous Access
Consider:
1__global__ void vector_add( 2 const float* A, 3 const float* B, 4 float* C, 5 int N 6) 7{ 8 int i = 9 blockIdx.x * blockDim.x 10 + threadIdx.x; 11 12 if (i < N) 13 { 14 C[i] = A[i] + B[i]; 15 } 16}
Suppose:
1blockDim.x = 256
The first warp has:
1threadIdx.x = 0 ... 31
Therefore:
1Thread 0 → A[0] 2Thread 1 → A[1] 3Thread 2 → A[2] 4... 5Thread 31 → A[31]
and:
1Thread 0 → B[0] 2Thread 1 → B[1] 3... 4Thread 31 → B[31]
This is an ideal basic access pattern.
4. Complete Working Example
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4__global__ void vector_add( 5 const float* A, 6 const float* B, 7 float* C, 8 int N 9) 10{ 11 int i = 12 blockIdx.x * blockDim.x 13 + threadIdx.x; 14 15 if (i < N) 16 { 17 C[i] = A[i] + B[i]; 18 } 19} 20 21int main() 22{ 23 const int N = 1 << 20; 24 25 size_t bytes = 26 N * sizeof(float); 27 28 float* h_A = 29 new float[N]; 30 31 float* h_B = 32 new float[N]; 33 34 float* h_C = 35 new float[N]; 36 37 for (int i = 0; i < N; i++) 38 { 39 h_A[i] = i; 40 h_B[i] = i * 2.0f; 41 } 42 43 float* d_A; 44 float* d_B; 45 float* d_C; 46 47 cudaMalloc(&d_A, bytes); 48 cudaMalloc(&d_B, bytes); 49 cudaMalloc(&d_C, bytes); 50 51 cudaMemcpy( 52 d_A, 53 h_A, 54 bytes, 55 cudaMemcpyHostToDevice 56 ); 57 58 cudaMemcpy( 59 d_B, 60 h_B, 61 bytes, 62 cudaMemcpyHostToDevice 63 ); 64 65 int threads = 256; 66 67 int blocks = 68 (N + threads - 1) 69 / threads; 70 71 vector_add<<<blocks, threads>>>( 72 d_A, 73 d_B, 74 d_C, 75 N 76 ); 77 78 cudaError_t err = 79 cudaGetLastError(); 80 81 if (err != cudaSuccess) 82 { 83 printf( 84 "Kernel launch error: %s\n", 85 cudaGetErrorString(err) 86 ); 87 88 return 1; 89 } 90 91 err = cudaDeviceSynchronize(); 92 93 if (err != cudaSuccess) 94 { 95 printf( 96 "Kernel execution error: %s\n", 97 cudaGetErrorString(err) 98 ); 99 100 return 1; 101 } 102 103 cudaMemcpy( 104 h_C, 105 d_C, 106 bytes, 107 cudaMemcpyDeviceToHost 108 ); 109 110 printf( 111 "C[0] = %.1f\n", 112 h_C[0] 113 ); 114 115 printf( 116 "C[100] = %.1f\n", 117 h_C[100] 118 ); 119 120 cudaFree(d_A); 121 cudaFree(d_B); 122 cudaFree(d_C); 123 124 delete[] h_A; 125 delete[] h_B; 126 delete[] h_C; 127 128 return 0; 129}
Expected:
1C[0] = 0.0 2C[100] = 300.0
5. Thread-to-Memory Mapping
This is one of the most important skills in CUDA.
For:
1int i = 2 blockIdx.x * blockDim.x 3 + threadIdx.x;
we get:
1Thread 0 → element 0 2Thread 1 → element 1 3Thread 2 → element 2 4...
So:
1Thread ID 2 ↓ 3Memory index 4 ↓ 5Physical memory address
When optimizing CUDA, always be able to draw this mapping.
6. What Is a Strided Access?
A strided access occurs when threads access memory with a gap.
For example:
1C[i * 4]
Then:
1Thread 0 → C[0] 2Thread 1 → C[4] 3Thread 2 → C[8] 4Thread 3 → C[12] 5...
Visualize:
1Thread Address 2 3T0 0 4T1 4 5T2 8 6T3 12 7T4 16
This is a stride of 4 elements.
Depending on the situation, this can reduce memory efficiency compared with contiguous access.
7. Stride Formula
A general strided access looks like:
1index = thread_id * stride;
For example:
1int i = 2 threadIdx.x * 4;
gives:
1i = 0 2i = 4 3i = 8 4i = 12 5...
The formula is:
1address = base + thread_id × stride
This formula is extremely useful for analyzing memory access patterns.
8. Contiguous vs Strided
Contiguous
1data[threadIdx.x]
Produces:
1T0 → 0 2T1 → 1 3T2 → 2 4T3 → 3
Stride 2
1data[threadIdx.x * 2]
Produces:
1T0 → 0 2T1 → 2 3T2 → 4 4T3 → 6
Stride 4
1data[threadIdx.x * 4]
Produces:
1T0 → 0 2T1 → 4 3T2 → 8 4T3 → 12
Large stride
1data[threadIdx.x * 1024]
Produces:
1T0 → 0 2T1 → 1024 3T2 → 2048 4T3 → 3072
The accesses become increasingly spread out.
9. Why Data Layout Matters
Consider a matrix:
1A = 2[ 1 2 3 4 ] 3[ 5 6 7 8 ] 4[ 9 10 11 12 ]
In row-major storage, it is laid out as:
11 2 3 4 5 6 7 8 9 10 11 12
If neighboring threads process neighboring columns of the same row:
1T0 → 1 2T1 → 2 3T2 → 3 4T3 → 4
excellent contiguous access.
10. Matrix Row Access
For:
1int row = ...; 2int col = ...; 3 4float value = 5 A[row * N + col];
If:
1T0 → col 0 2T1 → col 1 3T2 → col 2 4T3 → col 3
then:
1T0 → A[row*N + 0] 2T1 → A[row*N + 1] 3T2 → A[row*N + 2] 4T3 → A[row*N + 3]
This is contiguous.
11. Matrix Column Access
Now suppose:
1T0 → row 0, column 0 2T1 → row 1, column 0 3T2 → row 2, column 0 4T3 → row 3, column 0
Memory accesses become:
1A[0*N + 0] 2A[1*N + 0] 3A[2*N + 0] 4A[3*N + 0]
The difference between addresses is:
1N
So the access has a large stride.
For a large matrix, this can be significantly less memory-efficient than accessing adjacent elements.
12. Important Matrix Insight
For row-major matrices:
1A[row][col]
neighboring columns are adjacent in memory.
Therefore:
1Threads → columns
is often a favorable basic pattern.
Whereas:
1Threads → rows
for a fixed column can create a large stride.
This is one reason matrix kernels carefully arrange thread mappings.
13. 2D Thread Mapping
For matrix operations:
1int row = 2 blockIdx.y * blockDim.y 3 + threadIdx.y; 4 5int col = 6 blockIdx.x * blockDim.x 7 + threadIdx.x;
Then:
1C[row * N + col]
Each thread computes:
1Thread 2 ↓ 3(row, col) 4 ↓ 5C[row][col]
For a warp, the exact mapping of lanes to (row, col) matters enormously.
14. Complete Matrix Addition Kernel
1__global__ void matrix_add( 2 const float* A, 3 const float* B, 4 float* C, 5 int N 6) 7{ 8 int row = 9 blockIdx.y * blockDim.y 10 + threadIdx.y; 11 12 int col = 13 blockIdx.x * blockDim.x 14 + threadIdx.x; 15 16 if (row < N && col < N) 17 { 18 int index = 19 row * N + col; 20 21 C[index] = 22 A[index] + B[index]; 23 } 24}
Launch:
1dim3 threads(16, 16); 2 3dim3 blocks( 4 (N + 15) / 16, 5 (N + 15) / 16 6); 7 8matrix_add<<<blocks, threads>>>( 9 d_A, 10 d_B, 11 d_C, 12 N 13);
15. Why 16 × 16?
Because:
116 × 16 = 256 threads
and:
1256 / 32 = 8 warps
So one block contains:
18 warps
The block shape also creates a particular mapping between warp lanes and matrix coordinates.
But 16×16 is not universally optimal.
You must benchmark on the target GPU.
16. Warp Memory Access
Remember:
11 warp = 32 threads
For a simple 1D kernel:
1Warp 2├── Thread 0 3├── Thread 1 4├── ... 5└── Thread 31
Usually ask:
What addresses do these 32 threads access?
For:
1data[i]
you get:
1T0 → data[0] 2T1 → data[1] 3... 4T31 → data[31]
Great starting point.
17. The Four Questions You Should Ask
Whenever you see a CUDA memory access, ask:
Question 1
1Which thread accesses what element?
Question 2
1Are neighboring threads accessing neighboring elements?
Question 3
1Is the access aligned and efficient?
Question 4
1Is the data reused?
These four questions are extremely useful when reviewing kernels.
18. Memory Coalescing Is Not Just "Sequential"
A common beginner misunderstanding is:
"If threads access sequential memory, it's always perfectly optimized."
Not necessarily.
You also need to consider:
1data type 2alignment 3warp mapping 4cache behavior 5transaction size 6architecture 7access direction
The correct mental model is:
Coalescing is about how a warp's memory requests combine into efficient memory transactions.
19. Alignment
Suppose you're accessing:
1float
and each float is:
14 bytes
Then:
1data[0] → byte 0 2data[1] → byte 4 3data[2] → byte 8 4...
Contiguous threads access:
10 24 38 412 5...
This forms a contiguous region.
Alignment can affect how efficiently the memory system services the requests.
20. Vectorized Access
Sometimes data can be loaded in wider units.
For example:
1float4
contains:
14 floats
or:
116 bytes
Conceptually:
1float4 2├── x 3├── y 4├── z 5└── w
This can be useful in certain optimized memory-transfer patterns.
But don't automatically use float4 everywhere.
The compiler, architecture, alignment, and actual kernel bottleneck all matter.
21. Reading the Same Value
Now consider:
1Thread 0 → A[0] 2Thread 1 → A[0] 3Thread 2 → A[0] 4Thread 3 → A[0]
This isn't necessarily bad.
Cache mechanisms and special memory access behavior can make repeated reads efficient.
For shared memory, a broadcast-style access can also be handled efficiently in appropriate cases.
The important lesson:
Don't judge performance from one rule alone. Analyze the complete access pattern.
22. Transpose Example
Matrix transpose is a classic memory-access problem.
Input:
1A[row][col]
Output:
1B[col][row]
A naive transpose can have:
1Reads → coalesced 2Writes → strided
or the reverse depending on the mapping.
This is why optimized transpose kernels often use:
1Global Memory 2 ↓ 3Shared Memory Tile 4 ↓ 5Transpose in shared memory 6 ↓ 7Global Memory
Conceptually:
1Global 2 ┌──────────────┐ 3 │ A A A A │ 4 │ A A A A │ 5 │ A A A A │ 6 │ A A A A │ 7 └──────────────┘ 8 ↓ 9 Shared Memory 10 ↓ 11 transpose 12 ↓ 13 Global
This is a classic example of using shared memory to transform an inefficient global-memory access pattern into an efficient one.
23. Practical Transpose Kernel
Here's a learning implementation:
1#define TILE 32 2 3__global__ void transpose( 4 const float* input, 5 float* output, 6 int N 7) 8{ 9 __shared__ float tile[TILE][TILE]; 10 11 int x = 12 blockIdx.x * TILE 13 + threadIdx.x; 14 15 int y = 16 blockIdx.y * TILE 17 + threadIdx.y; 18 19 if (x < N && y < N) 20 { 21 tile[threadIdx.y][threadIdx.x] = 22 input[y * N + x]; 23 } 24 25 __syncthreads(); 26 27 int tx = 28 blockIdx.y * TILE 29 + threadIdx.x; 30 31 int ty = 32 blockIdx.x * TILE 33 + threadIdx.y; 34 35 if (tx < N && ty < N) 36 { 37 output[ty * N + tx] = 38 tile[threadIdx.x][threadIdx.y]; 39 } 40}
This demonstrates:
1Coalesced global read 2 ↓ 3Shared memory tile 4 ↓ 5Synchronization 6 ↓ 7Transpose 8 ↓ 9Coalesced global write
There are additional optimizations for avoiding shared-memory bank conflicts, which we'll study separately.
24. Why This Is Important for AI
This isn't only an academic CUDA concept.
Modern AI workloads constantly manipulate:
1[B, S, H] 2[B, H, S] 3[B, heads, sequence, head_dim]
and perform operations such as:
1transpose 2reshape 3permutation 4gather 5scatter 6matrix multiplication 7attention 8normalization
The physical memory layout can dramatically affect performance.
For example:
1Q 2[K] 3[V]
may be logically viewed one way but physically stored another way.
An optimized kernel needs to understand the actual layout.
25. Example: Transformer Tensor
Suppose:
1Q shape = 2[B, H, S, D]
where:
1B = batch 2H = attention heads 3S = sequence length 4D = head dimension
An attention kernel might access:
1Q[b][h][s][d]
The fastest mapping isn't simply:
1threadIdx.x → s
without thinking.
You need to determine:
1Which dimension is contiguous? 2Which dimension is reused? 3Which dimension should threads traverse? 4Which values fit in registers? 5Which data should be tiled?
This is the beginning of real GPU kernel engineering for AI.
26. Memory Access Optimization Process
When optimizing a kernel, follow this process:
11. Understand tensor shape 2 ↓ 32. Understand physical layout 4 ↓ 53. Map threads → elements 6 ↓ 74. Analyze one warp 8 ↓ 95. Check global memory access 10 ↓ 116. Check reuse 12 ↓ 137. Consider shared memory 14 ↓ 158. Check bank conflicts 16 ↓ 179. Check register usage 18 ↓ 1910. Benchmark
Don't optimize based purely on intuition.
Measure.
27. A Very Important Example
Suppose:
1int i = 2 blockIdx.x * blockDim.x 3 + threadIdx.x; 4 5C[i] = 6 A[i] + B[i];
Memory mapping:
1Warp 0 2 3T0 → A[0], B[0], C[0] 4T1 → A[1], B[1], C[1] 5T2 → A[2], B[2], C[2] 6... 7T31 → A[31], B[31], C[31]
This is a classic efficient pattern.
Now:
1int i = 2 (blockIdx.x * blockDim.x 3 + threadIdx.x) * 1024; 4 5C[i] = 6 A[i] + B[i];
Now:
1T0 → 0 2T1 → 1024 3T2 → 2048 4T3 → 3072 5...
This is a very different memory pattern.
You should immediately recognize:
Large stride → investigate memory efficiency.
28. Your Mental Model
You should now visualize CUDA memory access as:
1 WARP 2 │ 3 ┌───────┼────────┐ 4 ↓ ↓ ↓ 5 T0 T1 T2 6 │ │ │ 7 ↓ ↓ ↓ 8 Address Address Address 9 │ │ │ 10 └───────┼────────┘ 11 ↓ 12 Memory Transactions 13 ↓ 14 Cache / VRAM
Your job as a kernel programmer is to make this path efficient.
29. What You Should Be Able to Do Now
After this topic, you should be able to look at:
1int idx = 2 blockIdx.x * blockDim.x 3 + threadIdx.x; 4 5output[idx] = 6 input[idx];
and explain:
11. Thread calculates its index. 22. Threads in a warp get consecutive indices. 33. Consecutive threads access consecutive elements. 44. The global-memory access is coalescing-friendly. 55. Memory transactions can be serviced efficiently.
You should also be able to identify:
1input[threadIdx.x * 1024]
as a potentially problematic strided access.
30. Your CUDA Roadmap Progress
You've now covered:
1Phase 1 — CUDA Fundamentals 2├── GPU architecture ✅ 3├── CUDA kernels ✅ 4└── Thread indexing ✅ 5 6Phase 2 — GPU Parallelism 7├── Threads / Blocks / Grids ✅ 8├── Warps & SIMT ✅ 9├── Synchronization ✅ 10├── Shared Memory ✅ 11└── Memory Coalescing ✅ ← CURRENT
Next:
1Phase 3 — Advanced Memory Optimization 2│ 3├── Shared Memory Bank Conflicts 4├── Memory Transactions 5├── Cache Behavior 6├── Memory Bandwidth 7├── Pinned Memory 8├── Unified Memory 9└── Asynchronous Memory Transfers
Then the very important kernel-performance section:
1Phase 4 — Kernel Optimization 2│ 3├── Occupancy 4├── Register Pressure 5├── Warp Divergence 6├── Instruction-Level Parallelism 7├── Latency Hiding 8├── Kernel Fusion 9├── Tiling 10├── Reduction 11└── Profiling with Nsight
And finally for AI/LLM kernel programming:
1Phase 5 — AI GPU Kernels 2│ 3├── GEMM 4├── Softmax 5├── LayerNorm / RMSNorm 6├── Attention 7├── FlashAttention concepts 8├── Fused kernels 9├── Tensor Cores 10├── WMMA 11├── CUTLASS 12├── CUDA Graphs 13└── Custom PyTorch CUDA Extensions
The next topic I recommend is Shared Memory Bank Conflicts, because you now understand shared memory and coalescing; bank conflicts are the next major memory-performance problem you need to master before moving into occupancy and serious kernel optimization.