cuda - 使用张量核心时未注册共享内存负载

Question

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

__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 负载。