使用Nsight Compute无法检测RTX4070 Laptop Tensor Core占用率的问题

问题背景 #

• 硬件：RTX4070 Laptop（Ada Lovelace架构，SM 8.9）
• 工具：Nsight Compute（NCU）2023.3.0
• 现象：
  • 程序中手动实现了m8n8k4 Tensor Core MMA操作，同时调用了cusparseSpMV
  • NCU报告无活跃Tensor Core，Roofline图异常
  • 单独指定sm__inst_executed_pipe_tensor_op_*系列指标时，结果显示为N/A或0

自定义MMA内联函数代码 #

__forceinline__ __device__ void mma_m8n8k4_fp16_v2(half *a, half *b, float *c) {
    asm volatile (
        "mma.sync.aligned.m8n8k4.row.col.f16.f16.f32.f32 "
        "%0, %1, %2, %0;"
        : "+r"(c[0])
        : "r"(a[0]), "r"(b[0])
    );
}

排查原因 #

1. 指令架构兼容性问题
   RTX4070 Laptop属于Ada Lovelace架构，默认使用GMMA（第三代Tensor Core）指令，而m8n8k4是Volta/Turing时代的HMMA指令格式。旧指令在新架构上可能未被硬件优先调度，甚至被编译器转译为普通算术指令，导致Tensor Core未实际激活。

2. 编译优化导致指令被移除
   自定义内联汇编的MMA操作如果没有被有效调用，或者编译器判定其可以用更高效的普通指令替代，会直接优化掉Tensor Core指令，NCU自然无法检测到。

3. NCU指标适配错误
   sm__inst_executed_pipe_tensor_op_*系列指标是针对Volta/Turing架构的HMMA指令设计的，Ada架构的GMMA指令对应指标为sm__inst_executed_pipe_tensor_op_gmma_*，旧指标在新架构上会返回N/A或0。

4. cusparseSpMV未启用Tensor Core加速
   cusparseSpMV默认算法不一定支持Tensor Core，需要手动指定Tensor Core专用算法（如CUSPARSE_MV_ALG_MERGE_PATH_TENSOR），同时输入数据类型（FP16/TF32）、矩阵格式（如CSR对齐要求）也需符合条件，否则不会触发Tensor Core。

解决方法 #

• 替换为Ada兼容的GMMA指令
  将自定义的m8n8k4 HMMA指令替换为Ada架构支持的GMMA格式，例如m16n8k8的FP16指令，确保硬件能正确调度Tensor Core。示例指令：

  mma.sync.aligned.m16n8k8.row.col.f16.f16.f32.f32 %0, %1, %2, %0;
  

• 强制编译保留Tensor Core指令
  编译时添加-arch=sm_89 -ptxas-options=-v选项，查看编译输出是否包含Tensor Core指令（如gmma）；同时确保自定义MMA函数被实际调用，可通过添加输出验证逻辑避免被优化。

• 使用适配Ada的NCU指标
  更换统计指标为Ada架构专用的Tensor Core指标，例如：

  • sm__tensor_op_utilization：Tensor Core整体利用率
  • sm__inst_executed_pipe_tensor_op_gmma_f16：FP16类型GMMA指令执行数
  • sm__inst_executed_pipe_tensor_op_gmma_tf32：TF32类型GMMA指令执行数

• 配置cusparseSpMV启用Tensor Core
  调用cusparseSpMV前，通过cusparseSpMVAlgSetAttribute指定CUSPARSE_MV_ALG_MERGE_PATH_TENSOR算法；确保输入矩阵为FP16/TF32类型，CSR格式的行指针、列索引满足对齐要求；检查cusparseSpMV的返回值确认算法是否成功启用。

• 升级NCU版本
  升级到Nsight Compute 2024.x版本，新版本对Ada Lovelace架构的支持更完善，修复了旧版本的指标统计兼容性问题。

内容的提问来源于stack exchange，提问作者Severus Snape