Language: 简体中文 English

目录 (Table of Contents)

理解 GPU 性能

关键性能指标

指标 定义 目标值
GFLOPS 每秒十亿次浮点运算 接近理论峰值
内存带宽利用率 实际带宽 / 峰值带宽 > 80%
计算强度 FLOPs / Bytes 根据矩阵大小变化
Occupancy 活跃 warp / 最大 warp > 50%

Roofline 性能模型 

1
2
3
4
5
6
7
8
9

Performance (GFLOPS)
    ^
    |    ╭────────── 峰值计算性能
    |   /
    |  /  内存受限区域 | 计算受限区域
    | /
    |╱________________> 计算强度 (FLOPs/Byte)
         ↑
      Ridge Point

关键洞察：

  • 小矩阵：通常内存受限
  • 大矩阵：通常计算受限
  • 优化策略因矩阵大小而异

计算强度公式 

1
2
3
4
5
6
7
8

GEMM 计算强度 (Arithmetic Intensity):

AI = (2 × M × N × K) / ((M×K + K×N + M×N) × 4 bytes)

示例 (1024×1024×1024):
AI = 2 × 1024³ / (3 × 1024² × 4)
   = 2,147,483,648 / 12,582,912
   ≈ 170.7 FLOPs/Byte

Block Size 选择

基本参数

参数 说明 典型值范围
BM Block 处理的 M 维度 64 - 256
BN Block 处理的 N 维度 64 - 256
BK 每次迭代的 K 维度 8 - 32
TM 每线程处理的 M 维度 4 - 8
TN 每线程处理的 N 维度 4 - 8

针对不同矩阵大小的配置 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32

// 小矩阵 (< 512)
GemmConfig small = {
    .BLOCK_M = 64,
    .BLOCK_N = 64,
    .BLOCK_K = 8,
    .TM = 4,
    .TN = 4,
    .use_double_buffer = false,
    .use_vectorized_load = false
};

// 中等矩阵 (512 - 2048)
GemmConfig medium = {
    .BLOCK_M = 128,
    .BLOCK_N = 128,
    .BLOCK_K = 8,
    .TM = 8,
    .TN = 8,
    .use_double_buffer = true,
    .use_vectorized_load = true
};

// 大矩阵 (> 2048)
GemmConfig large = {
    .BLOCK_M = 128,
    .BLOCK_N = 256,
    .BLOCK_K = 16,
    .TM = 8,
    .TN = 8,
    .use_double_buffer = true,
    .use_vectorized_load = true
};

约束验证 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

═════════════════════════════════════════════════════════════
1. 线程数约束:
   threads_per_block = (BM / TM) × (BN / TN)
   threads_per_block ≤ 1024 (每个 block 的最大线程数)

2. 共享内存约束:
   shared_mem = (BM × BK + BK × BN) × sizeof(float)
                × (double_buffer ? 2 : 1)
   shared_mem ≤ 48KB (典型值) 或 96KB (A100)

3. 寄存器约束:
   registers_per_thread ≈ TM × TN + TM + TN + overhead
   registers_per_thread ≤ 255
   
   寄存器用量影响 Occupancy:
   - 每个 SM 有 65536 个寄存器
   - Occupancy = (block 数 × 线程数 × 每线程寄存器) / 65536
═════════════════════════════════════════════════════════════

AutoTuner 使用 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

#include "autotuner.h"

AutoTuner tuner;

// 自动调优并缓存结果
AutoTuner::Result result = tuner.tune(M, N, K, cublas_handle);

std::cout << "最优 Kernel: " 
          << kernel_type_name(result.config.kernel_type) << std::endl;
std::cout << "性能: " << result.gflops << " GFLOPS" << std::endl;

// 使用最优配置执行
tuner.execute_best(A, B, C, M, N, K, stream);

// 查看缓存的调优结果
tuner.print_cache();

内存优化技巧

1. 内存合并访问（Coalescing） 

1
2
3
4
5
6
7
8
9

// ✅ 好: 连续访问，合并内存事务
for (int i = 0; i < N; i += 4) {
    float4 val = load_float4(&data[threadIdx.x * 4 + i * blockDim.x * 4]);
}

// ❌ 差: 跨步访问，非合并
for (int i = 0; i < N; i++) {
    float val = data[threadIdx.x * N + i];  // 大跨步
}

2. Bank Conflict 避免 

1
2
3
4
5
6
7

// ❌ 有 Bank Conflict: 同 warp 线程访问同一 bank
__shared__ float smem[32][32];
float val = smem[threadIdx.x][0];  // 都在 bank 0

// ✅ 无 Bank Conflict: Padding 使访问分散到不同 bank
__shared__ float smem[32][33];  // +1 padding
float val = smem[threadIdx.x][0];  // 分散在 bank 0-31

Bank Conflict 模式：

访问模式 Bank Conflict 性能影响
smem[threadIdx.x][k] 无 1x
smem[k][threadIdx.x] 有（无 padding） 1/32x
smem[k][threadIdx.x] + padding 无 1x

3. 向量化加载 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

// 向量化加载函数
__device__ __forceinline__ float4 load_float4(const float* ptr) {
    return *reinterpret_cast<const float4*>(ptr);
}

// 使用向量化加载
void load_data_vectorized(float* dst, const float* src, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    
    // 每次处理 4 个元素
    for (int i = idx; i < n / 4; i += stride) {
        float4 val = load_float4(&src[i * 4]);
        dst[i * 4 + 0] = val.x;
        dst[i * 4 + 1] = val.y;
        dst[i * 4 + 2] = val.z;
        dst[i * 4 + 3] = val.w;
    }
}

4. 内存池优化 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

// 启用内存池统计
MemoryPool::instance().print_stats();

// 输出示例:
// MemoryPool Stats:
//   Active blocks: 10
//   Cached blocks: 5
//   Cache hits: 100
//   Cache misses: 20
//   Hit rate: 83.3%

// 根据命中率调整策略
if (hit_rate < 0.8) {
    // 增加预分配或调整块大小
}

GPU 架构特定优化

架构特性对比

架构 计算能力 关键特性 推荐配置
Volta 7.0 独立线程调度 BM=128, BN=128
Turing 7.5 Tensor Core 考虑 FP16
Ampere 8.0 异步拷贝 使用 cp.async
Ada 8.9 改进 Tensor Core 最大化计算密度
Hopper 9.0 Transformer Engine 使用 TMA

Volta (SM 7.0) 优化 

1
2
3
4
5
6
7
8
9
10

// Volta 特性: 独立线程调度
// 不需要显式同步即可安全使用 warp 原语

// Warp 级归约
__device__ float warp_reduce(float val) {
    for (int offset = 16; offset > 0; offset /= 2) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;
}

Ampere (SM 8.0+) 异步拷贝 

1
2
3
4
5
6
7
8

// Ampere 异步拷贝: 隐藏内存延迟
#if __CUDA_ARCH__ >= 800
__device__ void async_copy(float* smem, const float* gmem, int bytes) {
    __pipeline_memcpy_async(smem, gmem, bytes);
    __pipeline_commit();
    __pipeline_wait_prior(0);
}
#endif

FP16 Tensor Core 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

// 使用 WMMA API 访问 Tensor Core
#include <mma.h>
using namespace nvcuda::wmma;

__global__ void tensor_core_gemm(__half* A, __half* B, float* C, int M, int N, int K) {
    fragment<matrix_a, 16, 16, 16, __half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, __half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float> c_frag;
    
    fill_fragment(c_frag, 0.0f);
    
    for (int k = 0; k < K; k += 16) {
        load_matrix_sync(a_frag, A + k, K);
        load_matrix_sync(b_frag, B + k * N, N);
        mma_sync(c_frag, a_frag, b_frag, c_frag);
    }
    
    store_matrix_sync(C, c_frag, N, mem_row_major);
}

性能分析工具

NVIDIA Nsight Systems 

1
2
3
4
5

# 系统级分析
nsys profile -o profile_report ./benchmark

# 查看报告
nsys-ui profile_report.nsys-rep

NVIDIA Nsight Compute 

1
2
3
4
5
6
7
8
9
10
11
12

# Kernel 级详细分析
ncu --set full -o ncu_report ./benchmark

# 分析特定 kernel
ncu --kernel-name "optimized_gemm" ./benchmark

# 关键指标分析
ncu --metrics \
    sm__throughput.avg.pct_of_peak_sustained_elapsed," \
    memory__throughput.avg.pct_of_peak_sustained_elapsed," \
    launch__occupancy \
    ./benchmark

关键性能指标

指标 说明 健康值
Occupancy SM 利用率 > 50%
Memory Throughput 内存带宽利用率 > 80%
Compute Throughput 计算单元利用率 > 70%
L1/TEX Hit Rate L1 缓存命中率 > 90%
L2 Hit Rate L2 缓存命中率 > 80%

延迟分析 

1
2
3
4
5
6
7
8
9
10
11
12

# 分析 stall 原因
ncu --metrics \
    smsp__pcsamp_warps_issue_stalled_barrier," \
    smsp__pcsamp_warps_issue_stalled_long_scoreboard," \
    smsp__pcsamp_warps_issue_stalled_math_pipe_throttle," \
    smsp__pcsamp_warps_issue_stalled_membar," \
    smsp__pcsamp_warps_issue_stalled_mio_throttle," \
    smsp__pcsamp_warps_issue_stalled_no_instruction," \
    smsp__pcsamp_warps_issue_stalled_not_selected," \
    smsp__pcsamp_warps_issue_stalled_selected," \
    smsp__pcsamp_warps_issue_stalled_wait \
    ./benchmark

最佳实践

性能优化检查清单

  • 选择合适的 block size（参考矩阵大小建议）
  • 确保内存访问合并（使用 Nsight Compute 验证）
  • 避免 bank conflict（使用 padding）
  • 使用向量化加载（float4）
  • 启用 double buffering（大矩阵）
  • 检查 occupancy（目标 > 50%）
  • 分析 stall 原因
  • 考虑 kernel fusion（减少内存访问）
  • 使用 AutoTuner 自动调优
  • 在不同 GPU 上验证性能

开发流程建议 

1
2
3
4
5
6
7
8
9
10
11

1. 实现功能正确性
   ↓
2. 使用 Nsight Compute 分析瓶颈
   ↓
3. 针对性优化（内存/计算）
   ↓
4. 验证性能提升
   ↓
5. 回归测试（确保正确性）
   ↓
6. 文档记录优化策略

调试性能问题 

1
2
3
4
5
6
7
8
9
10
11
12
13
14

// 启用详细日志
Logger::instance().set_level(LogLevel::DEBUG);

// 使用 Profiler 分析
Profiler profiler;
auto result = profiler.profile("MyKernel", M, N, K, 5, 20, [&]() {
    my_kernel<<<grid, block>>>(...);
});

Profiler::print_result(result);

// 检查计算强度
float arithmetic_intensity = result.gflops / result.memory_bandwidth_gb;
std::cout << "计算强度: " << arithmetic_intensity << " FLOPs/Byte" << std::endl;

常见问题

Q: 为什么小矩阵性能差？

原因：

  • GPU 并行性无法充分利用
  • kernel 启动开销占比较大

解决方案：

  • 使用批量 GEMM（Batched GEMM）
  • 减小 block size
  • 考虑 CPU 执行（矩阵 < 128）
1
2

// 批量 GEMM
launch_batched_gemm_streams(desc, GemmKernelType::REGISTER_BLOCKED, handle);

Q: 为什么 double buffer 有时更慢？

原因：

  • Double buffer 增加共享内存使用
  • 可能导致 occupancy 下降

解决方案：

  • 检查共享内存使用量
  • 对于小矩阵不使用 double buffer
  • 使用 AutoTuner 自动选择

Q: 如何选择 FP16 vs FP32？

特性 FP16 FP32
内存带宽 2x 1x
Tensor Core ✅ ❌
精度 较低 标准
累加 FP32 FP32

建议：

  • 推理：使用 FP16（权重）+ FP32（累加）
  • 训练：使用 FP32
  • 对精度敏感：使用 FP32

Q: 如何检测 bank conflict？ 

1
2
3

# 使用 Nsight Compute
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld \
    ./benchmark

Q: Occupancy 低怎么办？

常见原因：

  • 寄存器用量过高
  • 共享内存用量过高
  • Block size 太小

解决方案：

  • 减少寄存器使用（例如减少临时变量）
  • 调整 block size
  • 减少共享内存使用
1
2
3

// 检查寄存器用量
__launch_bounds__(256, 4)  // 最大 256 线程，最小 4 个 block
__global__ void my_kernel(...) { ... }

相关链接

*最后更新：2025-04-16 文档版本：v1.1.0*

Back to top

MIT License | A learning project for the CUDA community