Repository navigation
V3. Vectorized and double buffer
(1) Vectorized memory access (float4): For data transfers between global memory and shared memory, or between shared memory and registers, use "float4" (128-bit) whenever possible to move 4 float values at once.
(2) Double buffering: Utilize two shared memory buffers (ping-pong). While computing the current tile using "buffer[cur]", prefetch the next tile from global memory into registers; once the computation is complete, write the prefetched data into "buffer[1-cur]".
Because the thread only processes the data block it is responsible for.
const float* Ab = A + (size_t)cRow * BM * K;
const float* Bb = B + cCol * BN;
float* Cb = C + (size_t)cRow * BM * N + cCol * BN;It is the same as V2.
"float4" represents the loading of 4 float values at once—128 bits in total—making it a 128-bit vector load.
the thread copies 4 floats each time.
For matrix A: 128 * 2, so 128 * 2 = 256 threads, 128 * 2 * 4 = 128 * 8 = BM * BK
For matrix B: 8 * 32, so 8 * 32 = 256 threads, 8 * 32 * 4 = 8 * 128 = Bk * BN
# each thread maps to 4 floats, 128 * 2 = 256.
const int innerRowA = threadIdx.x / (BK / 4); // 0..BM-1
const int innerColA = threadIdx.x % (BK / 4); // 0..(BK/4-1)
// for B, 8 * 32 = 256
const int innerRowB = threadIdx.x / (BN / 4); // 0..BK-1
const int innerColB = threadIdx.x % (BN / 4); // 0..(BN/4-1)full load shared memory one time:
256 * 4 = 8 * 128 = 1024
__shared__ float As[2][BK * BM];
For matrix A
matrix A make a transpose, A[m][k] -> As[k][m]
Facilitates reading by threads in the K direction.
As[0][(innerColA * 4 + 0) * BM + innerRowA] = tmpA.x;
As[0][(innerColA * 4 + 1) * BM + innerRowA] = tmpA.y;
As[0][(innerColA * 4 + 2) * BM + innerRowA] = tmpA.z;
As[0][(innerColA * 4 + 3) * BM + innerRowA] = tmpA.w;
# As[0][k * BM + m],
# so A in shared memory：
m →
k=0 A00 A10 A20 A30 ...
k=1 A01 A11 A21 A31 ...
k=2 A02 A12 A22 A32 ...
...For matrix B
//the thread maxtrix for B is (8 * 32)
float4 tmpB = reinterpret_cast<const float4*>( &Bb[(0 + innerRowB) * N + innerColB * 4] )[0];
reinterpret_cast<float4*>(&Bs[0][innerRowB * BN + innerColB * 4])[0] = tmpB;For matrix A:
forward 128 * 2, 4 float each time, (BM * BK)
For matrix B:
forward 8 * 32, 4 float each time, (BK * BN)
Loop:
after loop:
write result to global memory matrix C.
cd /content/gpu_perf
cmake -S . -B build -DCMAKE_CUDA_ARCHITECTURES=75 -DCMAKE_BUILD_TYPE=Release
cmake --build build -j
./build/v3_vecdbuf 512
-[verify] max_abs_err=9.536743e-06 -> PASS
# Measured Data
N | V2 (8x8) | V3 (8x8) | speedup v2
----| --------- | ----------| ----------
2048 | 2734 | 3203 | 1.17× sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --set full -o v3_profile ./build/v3_vecdbuf 2048
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --import v3_profile.ncu-rep --page details > v3_report.txt
grep -A 5 -i "bank conflict" v3_report.txt
grep -B 2 -A 8 "MIO" v3_report.txt | head -30conflict across all 16777216 shared load requests.This results in 33554432 bank conflicts, which represent
40.00% of the overall 83886080 wavefronts for shared loads. Check the Source Counters section for
conflict across all 2621440 shared store requests.This results in 2097152 bank conflicts, which represent
33.33% of the overall 6291456 wavefronts for shared stores. Check the Source Counters section for
uncoalesced shared stores.
...- Shared load (read) conflicts: 40.00% (33.5M conflicts out of 83.9M total wavefronts)—higher than the 33.33% rate seen with V2-4x4 padding.
- Shared store (write) conflicts: 33.33%, 2.4-way conflicts—this is a newly emerged issue; neither V2 nor V1 exhibited store conflicts.
To achieve high memory bandwidth, shared memory is physically organized into 32 memory banks of equal width that can be accessed simultaneously. Consecutive 128-byte segments of shared memory are distributed across the 32 banks at the same level; the bank is the smallest unit of shared memory.
Threads within the same block share a region of shared memory. A bank conflict occurs when multiple threads in a warp simultaneously access different addresses within the same bank; this causes the accesses to be serialized, thereby reducing performance.
Bank conflicts do not occur between threads in different warps.
[verify] PASS first.
Nsight Compute check the result.
- Shared load (read) conflicts eliminated.
- Shared store (write) conflicts still persist.
For size: 2048
[bench] avg_time=5.363 ms, GFLOPS= 3203.65 <- before fix
[bench] avg_time=4.461 ms, GFLOPS=
[verify] PASS first.
Nsight Compute check the result.
- Shared load (write) conflicts eliminated.
- The kernel is compute-bound, and its performance is already quite good. DRAM throughput utilization is 8%, the L2 cache hit rate is 80%, and there are no issues with branch divergence or register spills.
For size: 2048
[bench] 4.281 ms, GFLOPS=
faster than v2 1.25x
/content/gpu_perf# ncu --kernel-name regex:vectorized_dbuf_gemm --launch-count 1 \
--metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum,l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum,smsp__inst_executed_op_shared_ld.sum \
./build/v3_vecdbuf 2048
# result:
=== V3 Vectorized + Double-Buffered GEMM (BM=128 BN=128 BK=8 TM=8 TN=8) ===
M=2048 N=2048 K=2048
threads/block=256, each thread calculates 8 x 8 = 64 output elements
shared memory per block = 16640 bytes (limit 49152B, including double buffering x2)
==PROF== Connected to process 12337 (/content/gpu_perf/build/v3_vecdbuf)
[verify first with N<=512]
==PROF== Profiling "vectorized_dbuf_gemm_kernel": 0%....50%....100% - 1 pass
[bench] avg_time=8.006 ms, GFLOPS=2145.79
Note: Compare against V2 (TM=TN=8, 1974 GFLOPS @ N=2048)
==PROF== Disconnected from process 12337
[12337] v3_vecdbuf@127.0.0.1
vectorized_dbuf_gemm_kernel(const float *, const float *, float *, int, int, int) (16, 16, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
Section: Command line profiler metrics
-------------------------------------------------------- ----------- ------------
Metric Name Metric Unit Metric Value
-------------------------------------------------------- ----------- ------------
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum 14,366
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum 22,970
l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum 50,346,014
smsp__inst_executed_op_shared_ld.sum inst 16,777,216# Statically analyze the instruction composition of the entire kernel.
cuobjdump -sass ./build/v3_vecdbuf > v3.sass
# Count records by opcode
grep -E '^\s+/\*[0-9a-f]{4}\*/' v3.sass \
| sed -E 's/^\s+\/\*[0-9a-f]+\*\/\s+//; s/^@!?U?P[0-9T]+\s+//' \
| awk '{print $1}' | sed 's/\..*//' | sort | uniq -c | sort -rn | head -20
# Dynamic statistics (instructions actually executed) are more accurate:
ncu --kernel-name regex:vectorized_dbuf_gemm --launch-count 1 \
--metrics sass__inst_executed_per_opcode,smsp__sass_thread_inst_executed_op_ffma_pred_on.sum \
--print-metric-instances details \
./build/v3_vecdbufncu -i v3_full.ncu-rep --page raw --csv \
--metrics regex:smsp__average_warps_issue_stalled_.*_per_issue_active.ratio \
| tr -d '"' | awk -F, 'NR>1 && $NF+0>0.05 {print $(NF-0), $0}' | sort -rn | head -12they show that, at N=2048, the remaining gap is roughly divided into two parts: approximately 8.5% stems from the tail, and about 21% arises from stalls within the SM (primarily mio, barrier, lg, and short_scoreboard). The tail represents the largest single component that can currently be quantified.
V3 boosted GFLOPS to 1.25 times that of V2.
bench
