Repository navigation
V2. register blocking
if each thread computes a tile(4×4=16）, (64×64) / 16 =256 threads.
if each thread computes a tile(8×8=64）, (64×64) / 64 =64 threads.
# for example
float regM[TM] = {0.0f};
__global__ void kernel(...) { float x = 0.0f; }x is a thread-local variable, but the compiler might place it in a register.
How can you verify whether it actually ended up in a register?
Run # nvcc -Xptxas -v kernel.cu.
You will see output similar to: Used 32 registers, 0 bytes lmem.
is organized into multiple
If multiple threads simultaneously
The GPU must then serialize these accesses, resulting in reduced performance.
Each thread computes TM × TN elements (defaulting to 4×4=16).
In each iteration of the K-loop, only TM + TN values need to be loaded from shared memory into registers to perform TM × TN FMA operations.
This shifts the memory-to-compute ratio from 2:1 down to (TM + TN) : (TM × TN) = 8:16 or 0.5:1.
Within a block, the 256 threads are arranged in a threadRow (0–15) by threadCol (0–15) layout,
and a single warp (32 threads) encompasses two consecutive threadRow groups.
When reading As[threadRow*TM+i][k], the 16 threads within the same threadRow group
This constitutes a shared memory "broadcast" access rather than a bank conflict.
Consequently, this design is inherently free of bank conflicts and requires no additional padding techniques. The same logic applies to reading Bs (broadcast based on threadCol).
use registers for each iteration of the K-loop.
The figure below shows the movement of the pointer for matrix C.
# code in project
C += cRow * BM * N + cCol * BN;
The figure below illustrates the process of copying two data of A into shared memory.
# code in project
for (int loadOffset = 0; loadOffset < BM; loadOffset += strideA) {
As[(innerRowA + loadOffset) * BK + innerColA] =
A[(innerRowA + loadOffset) * K + innerColA];
}
The figure below illustrates the process of calculation of two registers.
# code in project
for (int resIdxM = 0; resIdxM < TM; ++resIdxM) {
#pragma unroll
for (int resIdxN = 0; resIdxN < TN; ++resIdxN) {
threadResults[resIdxM * TN + resIdxN] +=
regM[resIdxM] * regN[resIdxN];
}
}
# Measured Data
N | V1 | V2 (4x4) | speedup v1| V2 (8x8) | speedup v1
----| --- | -------- | ----------| -------- | ----------
2048 | 567 | 1399 | 2.47× | 1974 | 3.48× ./build/v2_reg8x8 2048
nvcc -O3 -arch=sm_75 -DTM=4 -DTN=4 --ptxas-options=-v src/v2_register_blocking.cu -o /tmp/v2_4x4_check
nvcc -O3 -arch=sm_75 -DTM=8 -DTN=8 --ptxas-options=-v src/v2_register_blocking.cu -o /tmp/v2_8x8_checkshow result:
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z29register_blocking_gemm_kernelPKfS0_Pfiii' for 'sm_75'
ptxas info : Function properties for _Z29register_blocking_gemm_kernelPKfS0_Pfiii
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 66 registers, used 1 barriers, 4096 bytes smem, 388 bytes cmem[0]
ptxas info : Compile time = 16.498 ms
-----------------------------------------
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z29register_blocking_gemm_kernelPKfS0_Pfiii' for 'sm_75'
ptxas info : Function properties for _Z29register_blocking_gemm_kernelPKfS0_Pfiii
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 121 registers, used 1 barriers, 4096 bytes smem, 388 bytes cmem[0]
ptxas info : Compile time = 60.577 msWhen TM=TN=4, each block has NUM_THREADS = (BM * BN) / (TM * TN) = 16 × 16 = 256.
When TM=TN=8, each block has NUM_THREADS = (BM * BN) / (TM * TN) = 8 × 8 = 64 threads.
The 8x8 configuration trades "fewer threads and higher per-thread register usage" for
"lower total register overhead per block" and "higher arithmetic intensity";
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --set full -o v2_profile ./build/v2_reg4x4 2048
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --import v2_profile.ncu-rep --page details > v2_report.txt
grep -A 5 -i "bank conflict" v2_report.txtshow result
conflict across all 134217728 shared load requests.This results in 134217728 bank conflicts, which
represent 50.00% of the overall 268435456 wavefronts for shared loads. Check the Source Counters section for
uncoalesced shared loads.
Section: Scheduler Statistics
---------------------------- ----------- ------------
...Increase the stored row width of As by 1 (changing BK=8 to BK+1=9). This makes the address interval across threadRow groups TM × 9 = 36; since 36 % 32 = 4, there is no longer a conflict within the same bank.
Load conflicts have indeed decreased:
Before fix: 268,435,456 total wavefronts, 50% conflicts
After fix: 201,326,592 total wavefronts, 33.33% conflicts (67,108,864 occurrences)
# Measured Data
N | V1 | V2 (4x4) | speedup v1| V2 (8x8) | speedup v1
----| --- | -------- | ----------| -------- | ----------
2048 | 567 | 1399 | 2.47× | 1974 | 3.48× <- Before padding
2048 | 567 | 1770 | 3.12× | 2734 | 4.82× <- After padding# ncu
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --set full -o v2_profile ./build/v2_reg8x8 2048
# save to v2_report.txt
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH"
/usr/local/cuda/bin/ncu --import v2_profile.ncu-rep --page details > v2_report.txt head -100 v2_report.txt
# get performance parameters
grep -A 12 "GPU Speed Of Light Throughput" v2_report.txt
# example output
Section: GPU Speed Of Light Throughput
----------------------- ----------- ------------
Metric Name Metric Unit Metric Value
----------------------- ----------- ------------
DRAM Frequency Ghz 5.00
SM Frequency Mhz 585.00
Elapsed Cycles cycle 7,198,177
Memory Throughput % 58.72
DRAM Throughput % 14.38
Duration ms 12.30
L1/TEX Cache Throughput % 95.83
L2 Cache Throughput % 9.91
SM Active Cycles cyclme 6,999,339.90
--
... Metric | V1 | V2(after padding) | change
---------------------- | ----------------- | ----------------- | ---------------------
Memory Throughput | 89.2% | 58.7% | ↓ Significant decline
DRAM Throughput | 9.6-9.8% | 13.7-14.8% | ↑ went up
L1/TEX Cache Throughput | ~89.8% | ~95.8% | ↑ Slight increase, approaching saturation.
L2 Cache Throughput | 4.5-4.6% | ~9.9% | ↑ Double
Duration（profiled） | ~53-55ms | ~12.3ms | Significantly faster
grep -A 15 "GPU Speed Of Light Throughput" v1_report.txt | grep -E "Compute \(SM\)|Memory Throughput"
grep -A 3 "Achieved Occupancy" v1_report.txt
grep -A 3 "Achieved Occupancy" v2_report.txt
grep -B 2 -A 8 "MIO" v2_report.txt | head -30
Metric | V1 | V2(after padding) | change
---------------------- | ----------------- | ----------------- | ---------------------
| Achieved Occupancy | 99.98%（full） | 52.5% | Cut in half
| Active Warps/SM | ~32 | ~16.8 | Cut in half
| Compute (SM) Throughput | 89.26% | 58.7% | decreasing rather than increasing
| MIO stall | 26.9 | 6.2-6.4（down 1/4）| Significant improvement
| cycles instructions | 39.89 | 11.6-11.7 | Dropped to one-thirdMemory throughput dropped from 89% to 59%,
indicating that register reuse indeed reduced the pressure on memory access instructions.
Register blocking cuts the absolute stall time caused by MIO operations to one-quarter (from 26.9 to 6.2 cycles),
but the trade-off is lower occupancy (dropping from 99.98% to 52.5%) due to increased register usage.
This is a trade-off, what truly matters is the actual, measured increase in GFLOPS.
Design assumption: “Reads are broadcast, so there should be zero conflicts.” → NCU measurements proved otherwise, revealing a 50% conflict rate → traced the issue to the classic pitfall where the stride is exactly a multiple of 32 → fixed it with padding → unexpectedly introduced write conflicts → after weighing the trade-offs, decided to keep the change (net effect: +14%).
Initial assumption: “The bottleneck should shift toward the compute units.” → measurements showed that both Compute Throughput and Memory Throughput actually decreased → instead of forcing the original hypothesis to fit the data, checked Achieved Occupancy → found it had been cut in half, from 99.98% to 52.5% → traced the real cause to the occupancy cost incurred by increased register pressure.