手写 CUDA 实现的通用矩阵乘法（GEMM），附带基准测试和性能分析工具。
| 变体 | 文件 | 说明 | N=2048 TFLOPS |
|---|---|---|---|
| CPU 基准 | gemm_cpu.cpp |
经典三重循环，O(N³) | ~0.001 |
| CUDA 朴素 | gemm_naive.cu |
每线程 1 个输出元素，全局内存 | ~1.03 |
| CUDA 分块 | gemm_tiled.cu |
共享内存分块（16×16），+1 padding | ~1.01 |
| CUDA 寄存器分块 | gemm_reg_blocked.cu |
寄存器分块（每线程 4×4），BM=BN=64 | ~6.88 |
| CUDA 寄存器分块 Vec4 | gemm_reg_blocked.cu |
+ float4 向量化协作加载 | ~7.11 |
| CUDA 双缓冲 8×8 | gemm_reg_blocked_v2.cu |
128×128 分块 + shared 双缓冲 + float4 写回 | ~7.0* |
| CUDA cp.async 8×8 | gemm_reg_blocked_v2.cu |
+ __pipeline_memcpy_async 异步拷贝 |
~8.9* |
| cuBLAS | gemm_cublas.cu |
NVIDIA 厂商优化的 cublasSgemm | ~11.55 |
在 RTX 4070 Laptop 上，reg_blocked_vec4 在 N=2048 达到约 7.11 TFLOPS，约为 cuBLAS（~11.55 TFLOPS）的 ~61%。主要提升来自寄存器分块提高算强；Vec4 作为进一步的内存访问优化带来较小但稳定的边际收益。
所有 GPU 基准测试仅计 内核执行时间（cudaEvent）。设备内存一次性预分配；H2D/D2H 传输、cudaMalloc 和 cuBLAS handle 创建均不计入。预热 10 次，重复 20 次取平均。
-
tiled 有时会略慢于 naive：本项目的 tiled 实现仍是"每线程计算 1 个 C 元素"，每个 K-tile 需要两次
__syncthreads()，共享内存读写与同步开销可能抵消了部分数据复用收益。因此在大矩阵（如 N=2048）上，tiled 可能与 naive 接近甚至略低，这是常见现象，并不代表实现错误。 -
reg_blocked 是主要性能跃迁点：寄存器分块（每线程 4×4 输出）显著提高 arithmetic intensity（每次从 shared 取到寄存器后做 16 次 FMA），通常能将 kernel 从偏 memory-bound 推向更 compute-heavy 的区域，因此 TFLOPS 会出现数量级提升。
-
Vec4 提升通常只有个位数 %：float4 向量化加载主要优化 global memory 的 load 指令与事务效率。当 reg_blocked 后 kernel 的主瓶颈更多来自计算吞吐、shared 访问、寄存器压力与调度（而非纯 global load），因此 Vec4 的边际收益会自然降低（例如 N=2048 提升 3–10% 属正常范围）。
关于
nvidia-smi显示 P8/低功耗：单次 GEMM kernel 往往只有 1–3ms（尤其 cuBLAS 更短），程序结束后 GPU 会立刻回到 P8 省电态，因此手动运行一次nvidia-smi很容易看到"空载"快照。若需观察高负载状态，请在程序中加入--stress <seconds>模式持续运行 5–10 秒，并使用nvidia-smi -l 1采样。
cmake -B build -DCMAKE_CUDA_ARCHITECTURES=89
cmake --build build --config Release（可选）也可使用 native 自动选择本机架构：
cmake -B build -DCMAKE_CUDA_ARCHITECTURES=native
cmake --build build --config Release依赖：CMake ≥ 3.18 · C++17 · CUDA Toolkit ≥ 11.0
# 所有模式，N=512，正确性检查
./build/gemm --mode all --n 512 --verify
# 单一模式，CSV 输出
./build/gemm --mode reg_blocked_vec4 --n 2048 --csv
# 模式：cpu | naive | tiled | reg_blocked | reg_blocked_vec4 | cublas | all
# 选项：--warmup <N> --repeat <N> --verify --csv可选：建议增加
--stress <seconds>（持续运行若干秒）用于观察 GPU 进入 P0/P2 状态与稳定采样（nvidia-smi -l 1）。该模式不影响 kernel-only 计时口径，仅用于压测/观察。
powershell benchmark/run_bench.ps1 # Windows
bash benchmark/run_bench.sh # Linux / WSL
python benchmark/plot.py # 生成图表输出：results/results.csv、results/performance_tflops.png、results/performance_time.png
| 步骤 | 技术 | 核心思路 |
|---|---|---|
| 1 | 朴素 | 每线程 1 个输出元素，全部读全局内存 |
| 2 | 共享内存分块 | 将 BK×TILE 大小的子块加载到 shared，减少全局读取 |
| 3 | 共享内存 padding | [TILE][TILE+1] 消除 bank conflict |
| 4 | 寄存器分块 | 每线程 4×4 个输出，每个 k 步 16 次 FMA |
| 5 | Vec4 加载 | 使用 float4 向量化加载改进 coalescing、减少 global load 指令（收益随瓶颈而变化） |
| 6 | 参数扩大 + 双缓冲 | BM=BN=128 TM=TN=8 每线程64 FMA + shared ping-pong 隐藏延迟 |
| 7 | cp.async 异步管线 | __pipeline_memcpy_async 绕过 L1 直接写入 shared（SM89 Ada 原生支持） |
| — | cuBLAS | NVIDIA 厂商优化 SGEMM 参考实现 |
所有内核通过 --verify 与 CPU 基准对比。最大误差 < 1e-3（FP32）。
ncu --set full -o results/profile build\Release\gemm.exe --mode reg_blocked_vec4 --n 2048
nsys profile -o results\timeline build\Release\gemm.exe --mode all --n 1024
cuda-gemm/
├── include/gemm.h # API 声明、常量定义
├── src/
│ ├── main.cpp # 命令行、仅内核基准测试驱动
│ ├── gemm_cpu.cpp # CPU 基准
│ ├── gemm_naive.cu # GPU 朴素内核
│ ├── gemm_tiled.cu # GPU 分块内核（共享 padding）
│ ├── gemm_reg_blocked.cu # GPU 寄存器分块 + Vec4 向量化内核
│ ├── gemm_reg_blocked_v2.cu # GPU 双缓冲 + cp.async 优化内核
│ ├── gemm_cublas.cu # cuBLAS 包装
│ └── utils.h # 矩阵初始化、误差计算
├── benchmark/
│ ├── run_bench.ps1 / .sh # 基准测试脚本
│ └── plot.py # 图表生成
├── results/ # 生成的输出
├── CMakeLists.txt
└── README.md