问题描述
我正在尝试在采用图灵架构设计的 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 (将#修改为@)