如何通过Nsight Compute查看我的张量核心占用率和利用率？

Question

在我的cuda程序中，我使用了许多张量核心运算，例如m8n8k4，甚至使用cusparseSpMV。然而，当检查ncu报告时，它显示如下： 我的程序中没有活动张量。屋顶线也显得很奇怪。我不知道这是怎么回事以及如何解决。 我的 GPU 是 RTX4070 笔记本电脑。 NCU版本是

NVIDIA (R) Nsight 计算命令行分析器 版权所有 (c) 2018-2023 NVIDIA 公司 版本 2023.3.0.0（内部版本 33266684）（公开发布）

我的命令行是

ncu -f --launch-count 50 --set full --export profile_rep ./my_program

我真的很困惑这个问题。感谢您抽出时间来帮助我！

我尝试使用命令行仅检查有关张量核心的指标。

ncu -f --launch-count 50 --metrics sm__inst_executed_pipe_tensor_op_hmma.sum,sm__inst_executed_pipe_tensor_op_dmma.sum,sm__inst_executed_pipe_tensor_op.sum,sm__inst_executed_pipe_tensor_op_hmma.avg.pct_of_peak_sustained_active --export profile_rep ./my_program

但是，指标显示 N/A 或 0。 我想知道为什么ncu无法显示张量核心的占用情况以及如何解决这个问题。 如果是由于我的错误，请随时指出并给我您的建议。

在我的内核中，我使用了这样的张量核心操作：

__device__ __forceinline__ void mma_m8n8k4_fp16_v2(half *acc, uint32_t *A, half *frag_b)
{
    uint32_t const *B = reinterpret_cast<uint32_t const *>(&frag_b[0]);
    uint32_t *C = reinterpret_cast<uint32_t *>(&acc[0]);

    asm volatile(
        "mma.sync.aligned.m8n8k4.row.col.f16.f16.f16.f16"
        " { %0, %1, %2, %3 }, "
        " { %4, %5 }, "
        " { %6, %7 }, "
        " { %0, %1, %2, %3 };"
        : "+r"(C[0]), "+r"(C[1]), "+r"(C[2]), "+r"(C[3]):
        "r"(A[0]), "r"(A[1]), "r"(B[0]), "r"(B[1])
    ); 
}

Answer 1

在这种情况下，nsightcompute并没有对你撒谎（我不这么认为。）根据你所展示的tensorcore（TC）操作，我不希望在该操作上运行时该操作的SASS代码中有任何实际的TC指令cc8.9 GPU，例如 RTX 40 系列 GPU。

这就是 NCU 报告的最直接原因。

为了证实这一点，我们可以使用 godbolt 研究每种情况下（cc8.9 和 cc7.0）生成的 SASS。我们将稍微修改您的代码以删除

__forceinline__

指令并添加必要的包含。在cc8.9案例中，生成的SASS没有显示TC操作的证据。 在cc7.0案例中，我们看到了预期的TC操作。 （您无法通过尝试在 cc8.9 设备上运行为 cc7.0 编译的代码来修改此处的行为；这不起作用。至少，根据编译设置，它将即时进行 JIT 重新编译。）

我相信这个“原因”可以在PTX指南中找到：

注意

mma.sync.m8n8k4 针对目标架构 sm_70 进行了优化，在其他目标架构上可能会大幅降低性能。

这并没有确切地说它将导致非 TC 操作，但这是我的猜想。 如果我必须猜测的话，我猜测无论出于何种原因，芯片设计者认为这个特定的 mma 操作不值得在未来的 TC 硬件中继续发扬光大。 因此在（至少部分）后续架构中，不存在TC路径，因此提供替代的非TC路径用于PTX级别的兼容性（这不是二进制兼容性的示例，但CUDA模型不需要二进制兼容性从较低的计算能力到较高的计算能力。）