使用 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 负载

解决方法

暂无找到可以解决该程序问题的有效方法,小编努力寻找整理中!

如果你已经找到好的解决方法,欢迎将解决方案带上本链接一起发送给小编。

小编邮箱:dio#foxmail.com (将#修改为@)