FAQ - 常见问题解答

本文档收集了 CUDA GEMM 优化和项目使用的常见问题。

---

环境与构建

Q1: 如何验证 CUDA 环境是否正确安装？

# 检查 nvcc 版本
nvcc --version

# 检查 GPU 驱动
nvidia-smi

# 检查 CUDA 路径
echo $CUDA_HOME

Q2: 构建时找不到 CUDA？

确保 CMAKE_CUDA_COMPILER 指向正确的 nvcc：

cmake --preset default -DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc

Q3: 测试被跳过？

GPU 测试需要可用的 NVIDIA GPU。如果没有 GPU，测试会自动跳过：

MINI_INFERENCE_REQUIRE_CUDA_DEVICE();  // 自动跳过无 GPU 环境

---

CUDA 编程

Q4: 什么是线程束 (warp)？

线程束是 NVIDIA GPU 的基本执行单元，包含 32 个线程。同一线程束内的线程以 SIMD 方式执行。

关键点：

• 线程束内的线程应尽量执行相同的代码路径（避免分支发散）
• 线程束同步是隐式的，不需要显式调用 __syncthreads()

Q5: 什么是 bank conflict？

共享内存被划分为 32 个 bank，每个 bank 宽度为 4 字节。当同一线程束内的多个线程访问同一 bank 的不同地址时，会发生 bank conflict，导致访问串行化。

避免方法：

// 可能产生 bank conflict
__shared__ float data[BLOCK_SIZE][BLOCK_SIZE];
float val = data[threadIdx.y][threadIdx.x];

// 添加 padding 避免 bank conflict
__shared__ float data[BLOCK_SIZE][BLOCK_SIZE + 1];

Q6: 什么是合并访存 (coalesced access)？

当同一线程束内的线程访问连续的全局内存地址时，这些访问可以被合并为一次内存事务，大幅提高吞吐量。

最佳实践：

// 合并访存：线程访问连续地址
float val = data[threadIdx.x + blockIdx.x * blockDim.x];

// 非合并访存：线程访问跨度地址
float val = data[threadIdx.y * width + threadIdx.x];  // 可能不合并

---

GEMM 优化

Q7: 为什么 Naive GEMM 性能这么差？

Naive GEMM 的主要瓶颈：

1. 内存访问重复：每个输出元素需要读取 K 个 A 元素和 K 个 B 元素
2. 无合并访存：线程访问模式不连续
3. 低计算强度：计算/访存比很低

Q8: Tiled GEMM 为什么能提升性能？

Tiled GEMM 使用共享内存缓存数据：

1. 减少全局内存访问：数据在共享内存中复用
2. 提高计算强度：每个加载的数据参与多次计算
3. 可预测的访问模式：更容易实现合并访存

Q9: 双缓冲如何隐藏延迟？

双缓冲使用两份共享内存缓冲区：

// 伪代码
while (more_tiles) {
    // 计算当前 tile，同时加载下一个 tile
    compute(buffer[current]);
    load(buffer[next]);
    swap(current, next);
}

计算和内存访问重叠，隐藏内存延迟。

Q10: 寄存器分块为什么是最重要的优化？

寄存器分块让每个线程计算一个小 tile：

1. 最大化寄存器使用：寄存器是最快的存储
2. 提高计算强度：每个加载的数据参与更多计算
3. 减少同步：线程内计算不需要同步

Q11: 为什么不直接使用 cuBLAS？

本项目是教学目的，目标是理解优化原理。生产环境应该使用 cuBLAS。

但理解原理的价值：

• 面试时能解释 GPU 性能优化
• 遇到 cuBLAS 不支持的场景可以自己实现
• 理解 CUTLASS 等高级库的设计

---

性能分析

Q12: 如何使用 Nsight Compute？

# 详细分析单个 kernel
ncu --set full ./benchmark --kernel=vectorized

# 只看关键指标
ncu --metrics gpu__time_duration.sum ./benchmark

Q13: 如何解读性能数据？

关键指标：

指标	含义	目标
SM Efficiency	SM 利用率	> 80%
Memory Throughput	内存吞吐量	接近峰值
Warp Execution Efficiency	线程束执行效率	> 90%
Achieved Occupancy	实际占用率	取决于 kernel

Q14: 为什么我的 kernel 性能不稳定？

可能原因：

1. GPU 频率波动：检查 nvidia-smi -q -d CLOCK
2. 热节流：检查温度 nvidia-smi -q -d TEMPERATURE
3. 系统负载：关闭其他 GPU 应用
4. 内存碎片：重启程序或重启系统

---

项目使用

Q15: 如何添加新的 GEMM kernel？

1. 在 include/kernels.cuh 声明 kernel
2. 在 src/ 创建实现文件
3. 在 tests/ 添加测试
4. 在 benchmarks/ 添加 benchmark

Q16: 如何使用 AutoTuner？

#include "autotuner.h"

AutoTuner tuner;
tuner.add_param("BLOCK_SIZE", {16, 32, 64, 128});
tuner.add_param("TILE_M", {4, 8, 16});
auto best = tuner.search(gemm_kernel);

Q17: 如何使用 MemoryPool？

#include "memory_pool.h"

MemoryPool pool(1024 * 1024 * 100);  // 100 MB pool
auto* ptr = pool.allocate<float>(1024);
// 使用 ptr...
pool.deallocate(ptr);

---

理论问题

Q18: 什么是 Roofline 模型？

Roofline 模型是一个可视化性能分析框架：

Performance = min(Peak FLOPS, Peak BW * AI)

其中 AI (Arithmetic Intensity) = FLOPs / Bytes

意义：

• AI 低的 kernel 受内存带宽限制
• AI 高的 kernel 受计算能力限制

Q19: 如何计算 GEMM 的算术强度？

对于 M×N×K 的 GEMM：

FLOPs = 2 * M * N * K  (每个输出元素需要 K 次乘法和 K-1 次加法)
Bytes = 4 * (M*K + K*N + M*N)  (假设 float32)
AI = FLOPs / Bytes = 2*M*N*K / (4*(M*K + K*N + M*N))

对于方阵 (M=N=K)：

AI = 2*M^3 / (12*M^2) = M/6

Q20: 为什么 FP16 能提升性能？

1. 内存带宽减半：数据量减少
2. Tensor Core：利用专用硬件
3. 更多并行：相同寄存器数量下存储更多数据

---

更多问题？

如果你有其他问题，欢迎：

1. 在 GitHub Issues 提问
2. 查阅 CUDA Best Practices Guide
3. 阅读 项目文档