1. What Is a Memory Transaction?
A memory transaction is a unit of data transfer between a CUDA processor/cache and the memory system.
The simplified idea is:
1CUDA Threads 2 ↓ 3Memory Requests 4 ↓ 5Memory Transactions 6 ↓ 7Cache / Global Memory
A warp contains:
132 threads
Those 32 threads can generate memory requests.
The GPU's memory system analyzes those requests and services them using memory transactions.
2. Why Should You Care?
Suppose a warp needs:
1data[0] 2data[1] 3data[2] 4... 5data[31]
The addresses are close together.
The GPU can generally service these accesses efficiently.
But suppose the warp needs:
1data[0] 2data[1000] 3data[2000] 4data[3000] 5...
Now the addresses are scattered.
The memory system may need substantially more transactions.
Conceptually:
1Good access 2 332 threads 4 ↓ 5nearby addresses 6 ↓ 7few efficient transactions
versus:
1Poor access 2 332 threads 4 ↓ 5scattered addresses 6 ↓ 7many transactions 8 ↓ 9lower efficiency
3. Memory Request vs Memory Transaction
These are not exactly the same thing.
A thread generates a memory request.
For example:
1Thread 0 → data[0]
A warp can generate many requests:
1T0 → data[0] 2T1 → data[1] 3T2 → data[2] 4...
The memory system then combines/organizes those requests into appropriate memory transactions.
Conceptually:
1Thread requests 2 ↓ 3Coalescing 4 ↓ 5Memory transactions 6 ↓ 7Cache / DRAM
This is why coalescing matters.
4. Simple Example
Suppose each element is:
1float = 4 bytes
A warp accesses:
1T0 → data[0] 2T1 → data[1] 3T2 → data[2] 4... 5T31 → data[31]
The memory addresses are:
10 24 38 412 5... 6124
So the warp accesses a contiguous region:
1124 + 4 = 128 bytes
Conceptually:
1┌─────────────────────────────────────┐ 2│ 128 bytes │ 3│ 0 4 8 12 ... 124 │ 4└─────────────────────────────────────┘ 5 ↑ 6 Warp accesses
The hardware can service this efficiently.
Do not memorize "one warp always equals one 128-byte transaction." Actual transaction behavior depends on GPU architecture, cache level, alignment, access size, and other factors.
The important principle is:
A warp accessing a compact, well-aligned region is generally more efficient than a warp accessing scattered regions.
5. What Happens With Scattered Access?
Consider:
1float value = 2 data[threadIdx.x * 1024];
For the first warp:
1T0 → data[0] 2T1 → data[1024] 3T2 → data[2048] 4T3 → data[3072] 5...
Memory addresses:
10 24096 38192 412288 5...
The addresses are far apart.
Conceptually:
1T0 → █ 2T1 → █ 3T2 → █ 4T3 → █
The memory system cannot combine these requests as efficiently as a contiguous access.
6. Coalescing
Memory coalescing means that memory requests from threads in a warp can be serviced efficiently because the addresses have a favorable spatial relationship.
Good:
1data[threadIdx.x]
Bad example:
1data[threadIdx.x * 1024]
Potentially problematic:
1data[threadIdx.x * stride]
when stride is large.
7. Alignment
Alignment matters.
Suppose:
1float = 4 bytes
and you have:
1data[0] 2data[1] 3data[2] 4...
The accesses are naturally aligned when the base pointer is appropriately aligned.
But if your data starts at an awkward byte offset, the accesses can cross memory boundaries inefficiently.
Conceptually:
1Aligned: 2 3|-------- region --------| 40 127 5 ↑ 6contiguous warp access
versus:
1Misaligned: 2 3 |-------- region --------| 4 ↑ 5 access begins here
The same number of bytes may require additional memory-system work.
8. Why Data Type Matters
Consider:
1float
which is:
14 bytes
Now compare:
1double
which is:
18 bytes
and:
1half
which is:
12 bytes
A warp accessing:
1data[threadIdx.x]
covers a different number of bytes depending on the element type.
For example:
132 × 4 bytes = 128 bytes
for float.
For half:
132 × 2 bytes = 64 bytes
For double:
132 × 8 bytes = 256 bytes
The memory system and transaction behavior depend on the architecture and access size.
9. Memory Transaction Efficiency
A useful conceptual metric is:
1Requested Bytes 2──────────────────── 3Transferred Bytes
For example:
1Requested = 128 bytes 2Transferred = 128 bytes
Very efficient.
But if:
1Requested = 128 bytes 2Transferred = 512 bytes
then much more data was transferred than the kernel actually needed.
That indicates poor spatial efficiency.
10. Wasted Memory Traffic
Suppose:
1Thread 0 → 1 float 2Thread 1 → 1 float 3... 4Thread 31 → 1 float
The warp needs:
132 × 4 = 128 bytes
If the memory access pattern causes the GPU to fetch many separate regions, the actual memory traffic can be substantially larger.
Conceptually:
1Useful data 2 ↓ 3128 bytes 4 5Actual transferred 6 ↓ 7much larger
This is why a kernel can have:
1low FLOPs
but still be slow.
The bottleneck may be memory traffic.
11. Memory Transaction Example
Let's compare two kernels.
Kernel A
1__global__ void good( 2 const float* input, 3 float* output 4) 5{ 6 int i = 7 blockIdx.x * blockDim.x 8 + threadIdx.x; 9 10 output[i] = 11 input[i]; 12}
Mapping:
1T0 → input[0] 2T1 → input[1] 3T2 → input[2] 4...
Good.
Kernel B
1__global__ void bad( 2 const float* input, 3 float* output 4) 5{ 6 int i = 7 blockIdx.x * blockDim.x 8 + threadIdx.x; 9 10 output[i] = 11 input[i * 1024]; 12}
Mapping:
1T0 → input[0] 2T1 → input[1024] 3T2 → input[2048] 4...
Potentially poor memory efficiency.
12. Complete Benchmark Example
Let's build a simple experiment.
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4__global__ void good_kernel( 5 const float* input, 6 float* output, 7 int N 8) 9{ 10 int i = 11 blockIdx.x * blockDim.x 12 + threadIdx.x; 13 14 if (i < N) 15 { 16 output[i] = 17 input[i]; 18 } 19} 20 21__global__ void strided_kernel( 22 const float* input, 23 float* output, 24 int N, 25 int stride 26) 27{ 28 int i = 29 blockIdx.x * blockDim.x 30 + threadIdx.x; 31 32 int index = 33 i * stride; 34 35 if (index < N) 36 { 37 output[i] = 38 input[index]; 39 } 40} 41 42int main() 43{ 44 const int N = 1 << 24; 45 46 size_t bytes = 47 N * sizeof(float); 48 49 float* d_input; 50 float* d_output; 51 52 cudaMalloc( 53 &d_input, 54 bytes 55 ); 56 57 cudaMalloc( 58 &d_output, 59 bytes 60 ); 61 62 cudaMemset( 63 d_input, 64 0, 65 bytes 66 ); 67 68 int threads = 256; 69 70 int blocks = 71 (N + threads - 1) 72 / threads; 73 74 good_kernel<<<blocks, threads>>>( 75 d_input, 76 d_output, 77 N 78 ); 79 80 cudaDeviceSynchronize(); 81 82 cudaFree(d_input); 83 cudaFree(d_output); 84 85 return 0; 86}
This demonstrates the basic access pattern.
For a proper performance comparison, you should use CUDA events rather than CPU timing.
13. Correct CUDA Timing
For GPU kernels, use:
1cudaEvent_t start; 2cudaEvent_t stop; 3 4cudaEventCreate(&start); 5cudaEventCreate(&stop); 6 7cudaEventRecord(start); 8 9kernel<<<blocks, threads>>>( 10 ... 11); 12 13cudaEventRecord(stop); 14 15cudaEventSynchronize(stop); 16 17float milliseconds; 18 19cudaEventElapsedTime( 20 &milliseconds, 21 start, 22 stop 23); 24 25printf( 26 "Time: %f ms\n", 27 milliseconds 28);
Why?
Because CUDA kernel launches are generally asynchronous with respect to the CPU.
This:
1kernel<<<...>>>();
doesn't mean:
1CPU waits until kernel finishes
immediately.
CUDA events let you measure GPU execution much more appropriately.
14. Memory Bandwidth
Suppose your kernel reads:
1128 bytes
and writes:
1128 bytes
Total:
1256 bytes
If it takes:
11 microsecond
then:
1Bandwidth = 2256 bytes / 1 microsecond
which is:
1256 GB/s
because:
11 GB/s = 10^9 bytes/s
This is the basic idea behind effective memory bandwidth.
15. Effective Bandwidth Formula
Remember this formula:
1Effective Bandwidth = 2Total Bytes Transferred 3────────────────────── 4Execution Time
For a simple vector copy:
1Read: 2N × sizeof(float) 3 4Write: 5N × sizeof(float)
Total:
12 × N × sizeof(float)
Therefore:
1Bandwidth = 22 × N × sizeof(float) 3────────────────────── 4time
16. Example
Suppose:
1N = 1,000,000
and:
1float = 4 bytes
Read:
14 MB
Write:
14 MB
Total:
18 MB
If the kernel takes:
10.1 ms
then:
1Bandwidth = 28 MB / 0.1 ms
Approximately:
180 GB/s
depending on whether you use decimal or binary units.
17. Memory-Bound Kernel
Consider:
1C[i] = A[i] + B[i];
Each element requires approximately:
1Read A → 4 bytes 2Read B → 4 bytes 3Write C → 4 bytes
Total:
112 bytes
Computation:
11 addition
Very little arithmetic compared with memory traffic.
Therefore this type of operation is often memory-bandwidth bound.
Conceptually:
1Memory 2 ↓ 3████████████████ 4 ↓ 5 GPU
The arithmetic units may have plenty of unused capacity.
18. Compute-Bound Kernel
Now consider:
1Matrix Multiplication
Each loaded value can be reused many times.
Conceptually:
1Load data 2 ↓ 3Compute 4Compute 5Compute 6Compute 7Compute 8...
The arithmetic workload becomes much larger relative to memory traffic.
This can become:
1Compute-bound
19. Arithmetic Intensity
This connects directly to your previous lesson.
Formula:
1Arithmetic Intensity = 2Operations 3────────── 4Bytes moved
Low:
1few operations / many bytes
→ likely memory-bound.
High:
1many operations / relatively few bytes
→ potentially compute-bound.
This concept is central to GPU performance analysis.
20. Roofline Model
A very important concept in performance engineering is the Roofline Model.
It gives you a mental model:
1Performance 2 │ 3 │ Compute Roof 4 │ ─────────────────── 5 │ / 6 │ / 7 │ / 8 │ / 9 │ / 10 └────────────────────────────── 11 Arithmetic Intensity
There are two major limits:
1Memory bandwidth
and:
1Compute throughput
A kernel's achievable performance is limited by whichever resource is the bottleneck.
Conceptually:
1Performance ≤ min( 2 memory_bandwidth × arithmetic_intensity, 3 peak_compute 4)
This formula is extremely important.
21. Example of Roofline Thinking
Suppose a GPU has:
1Memory bandwidth = 500 GB/s 2Peak compute = 20 TFLOPS
Kernel arithmetic intensity:
110 FLOPs/byte
Memory-limited performance:
1500 GB/s × 10 FLOPs/byte 2= 35000 GFLOPS 4= 55 TFLOPS
Peak compute:
120 TFLOPS
Therefore:
1Potential performance limit 2≈ 5 TFLOPS
The kernel is likely memory-bound at that arithmetic intensity.
22. Why This Matters for Kernel Optimization
Suppose you optimize arithmetic:
110 operations 2→ 38 operations
But the kernel is memory-bound.
You may see little improvement.
Why?
Because:
1Memory 2 ↓ 3Bottleneck
Instead, you may need to optimize:
1coalescing 2memory traffic 3cache reuse 4shared-memory reuse 5data layout
This is why knowing the bottleneck is more important than randomly optimizing code.
23. Memory Transactions + Cache
The actual path can be more complex:
1Thread 2 ↓ 3Register 4 ↓ 5L1 Cache 6 ↓ 7L2 Cache 8 ↓ 9VRAM
If data is found in L1:
1Thread 2 ↓ 3L1 hit 4 ↓ 5data
If not:
1L1 miss 2 ↓ 3L2 4 ↓ 5possibly VRAM
Therefore, "global memory access" does not necessarily mean every request goes directly to physical VRAM.
24. Cache Hit
Suppose a thread accesses:
1data[100]
and later another thread accesses:
1data[100]
The data may still be available in cache.
Conceptually:
1First access 2 ↓ 3L1 miss 4 ↓ 5L2 / memory 6 ↓ 7cache 8 9Second access 10 ↓ 11L1/L2 hit 12 ↓ 13faster
This is why you should analyze both:
1memory access pattern
and:
1data reuse
25. Important Distinction
Don't think:
1Global Memory = Always Slow VRAM Access
Instead:
1Global memory address 2 ↓ 3Cache hierarchy 4 ↓ 5Possible cache hit 6 ↓ 7Possible memory access
Modern GPUs have sophisticated memory systems.
26. Memory Transactions in AI
Consider a tensor:
1X[B][S][H]
Suppose:
1B = 1 2S = 2048 3H = 4096
A kernel might process:
1X[b][s][h]
If neighboring threads access neighboring h values:
1T0 → X[0][s][0] 2T1 → X[0][s][1] 3T2 → X[0][s][2] 4...
you get a contiguous access pattern.
But if:
1T0 → X[0][0][h] 2T1 → X[0][1][h] 3T2 → X[0][2][h] 4...
then the stride may be:
1H
which could produce a much less efficient access pattern.
This matters heavily in:
1Attention 2LayerNorm 3RMSNorm 4GEMM 5Embedding 6Transpose 7Softmax
27. Real Kernel Optimization
Suppose your kernel is:
1Memory-bound
Your optimization path should be:
11. Check coalescing 2 ↓ 32. Check transaction efficiency 4 ↓ 53. Check cache hit rate 6 ↓ 74. Reduce unnecessary loads 8 ↓ 95. Increase data reuse 10 ↓ 116. Consider shared memory 12 ↓ 137. Consider kernel fusion
Not:
1Change random arithmetic
28. Kernel Fusion
Suppose you have:
1Kernel A 2 ↓ 3Global Memory 4 ↓ 5Kernel B 6 ↓ 7Global Memory
You might instead fuse them:
1Kernel A + B 2 ↓ 3Registers / Shared Memory 4 ↓ 5Global Memory
Example:
Without fusion:
1A → memory 2memory → B 3B → memory
With fusion:
1A 2↓ 3register 4↓ 5B 6↓ 7memory
This can dramatically reduce memory traffic when applicable.
This is one reason fused AI kernels are so important.
29. The Key Performance Equation
For kernel optimization, remember:
1Execution Time 2≈ 3max( 4 computation time, 5 memory time 6)
This is simplified, but useful.
If:
1memory time > computation time
you're likely memory-bound.
If:
1computation time > memory time
you're likely compute-bound.
Real kernels can have more complex bottlenecks and overlap, but this is a good starting model.
30. Your Current CUDA Mental Model
You've now built this:
1CUDA Program 2 │ 3 ↓ 4Kernel 5 │ 6 ↓ 7Grid 8 │ 9 ↓ 10Blocks 11 │ 12 ↓ 13Warps 14 │ 15 ↓ 16Threads 17 │ 18 ↓ 19Memory Access 20 │ 21 ┌───┴───────────────┐ 22 ↓ ↓ 23Global Memory Shared Memory 24 ↓ ↓ 25Coalescing Bank Conflicts 26 ↓ ↓ 27Transactions Banks 28 ↓ 29Cache 30 ↓ 31Memory Bandwidth
And above that:
1Arithmetic Intensity 2 ↓ 3Roofline 4 ↓ 5Memory-bound vs Compute-bound
This is the beginning of real GPU performance engineering.
31. Practical Exercise
Write three kernels:
Kernel 1 — Contiguous
1output[i] = input[i];
Kernel 2 — Stride 2
1output[i] = input[i * 2];
Kernel 3 — Large stride
1output[i] = input[i * 1024];
Then measure them with CUDA events.
Observe:
1Kernel 2 ↓ 3Execution time 4 ↓ 5Effective bandwidth
Then use NVIDIA profiling tools to investigate why they differ.
32. Phase 3 Progress
Your roadmap is now:
1Phase 3 — Advanced Memory Optimization 2 3├── Shared Memory Bank Conflicts ✅ 4├── Memory Transactions ✅ 5├── Cache Behavior ← NEXT 6├── Memory Bandwidth 7├── Pinned Memory 8├── Unified Memory 9└── Asynchronous Memory Transfers
Next: CUDA Cache Behavior
We'll cover:
1L1 cache 2L2 cache 3cache hits 4cache misses 5spatial locality 6temporal locality 7cache-friendly access 8cache-unfriendly access 9read-only data 10cache reuse 11and how to profile cache performance
After that, Memory Bandwidth will connect everything you've learned so far into actual performance calculations.