使用 Tensor Core 时未注册共享内存加载

Question

我正在尝试在采用图灵架构设计的 GPU 上使用 Tensor Core 来乘以 8x8 大小的块。为此，我使用 WMMA API 和大小为 16x16 的片段。我的假设是共享内存带宽会被浪费，因为加载到片段中的大多数数据并不代表有用的信息。在尝试量化时，我遇到了以下问题：Nsight Compute 上甚至没有报告使用

wmma::load_matrix_sync

的共享内存负载。为了测试这一点，我正在使用这个内核：

__global__
void test() {
    extern __shared__  half shmem[];
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;
    wmma::load_matrix_sync(a_frag, shmem, 16);
    wmma::load_matrix_sync(b_frag, shmem, 16);
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    wmma::store_matrix_sync((float*)shmem, c_frag, 16, wmma::mem_row_major);
}

Nsight Compute 报告共享内存存储，但不报告负载。这里发生了什么？我尝试了几种变体，但它仍然显示 0 负载。

Answer 1

根据评论中的信息添加答案：

新的 LDSM 指令未计入用于共享内存访问的 SM 硬件计数器中。 Nsight Compute 2020.3.1 中进行了修复。请参阅发行说明此处

正如假设的那样，在发布问题时，Nsight 计算不包括 wmma 生成的加载指令的指令计数。此问题已得到纠正。