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

如何解决使用 Tensor Core 时未注册共享内存负载

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

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

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