cuda - Tensor コアの使用時に共有メモリの読み込みが登録されない

Question

Turing アーキテクチャで設計された GPU で Tensor コアを使用して、サイズ 8x8 のブロックを乗算しようとしています。そのために、WMMA API とサイズ 16x16 のフラグメントを使用しています。私の推測では、フラグメントにロードされたほとんどのデータは有用な情報を表していないため、共有メモリの帯域幅が浪費されるというものでした。数値化しようとしているときに、次の問題に遭遇しました: wmma::load_matrix_sync を使用した共有メモリのロードは、Nsight コンピューティングでも報告されません。それをテストするために、私はこのカーネルを使用しています:

__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 と表示されます。