CUDA：如何在具有计算能力>= 7.2的设备上检测设备内共享内存银行冲突？

Question

在具有计算能力小于等于7.2的设备上，我总是使用 nvprof --events shared_st_bank_conflict 但当我在CUDA10上运行RTX2080ti时，它返回了 Warning: Skipping profiling on device 0 since profiling is not supported on devices with compute capability greater than 7.2 那么我该如何检测这些设备是否存在共享内存银行冲突呢？

我已经安装了Nvidia Nsight Systems和Nsight Compute，但没有找到这样的分析报告...

谢谢！

- LiMou

---

你应该使用Nsight Compute。 - Robert Crovella

Answer 1

你可以使用--metrics：

其中，

nv-nsight-cu-cli --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum

针对从共享内存中读取（载入）时可能发生的冲突，或者

nv-nsight-cu-cli --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum

在写入（存储）共享内存时可能会出现冲突。

- Deepak C U

Answer 2

看起来这是一个问题，并且在this NVIDIA论坛帖子中得到了解决。显然可以使用其中一个Nsight工具（CLI或UI）来支持它。

- user11127113

Answer 3

正如其他人指出的那样，nvprof已被Nsight Compute取代，请查看它们的度量等效映射。

特别是，`shared_efficiency`被映射为`smsp__sass_average_data_bytes_per_wavefront_mem_shared`（晦涩难懂！）。

我有一种感觉，在这个过渡期间，更多的度量指标受到了影响。开玩笑的话不说了，让我们来演示如何使用它。为此，我们来看一个故意引起冲突的内核。

__global__ void kernel(int offset)
{
    __shared__ unsigned int sharedMem[4096];

    int threadId = threadIdx.x;

    // init shared memory
    if (threadId == 0)
    {
        for (int i = 0; i < 4096; i++) sharedMem[i] = 0;
    }
    __syncthreads();

    // repeatedly read and write to shared memory
    unsigned int index = threadId * offset;
    for (int i = 0; i < 10000; i++)
    {
        sharedMem[index] += index * i;
        index += 32;
        index %= 4096;
    }
}

这个内核应该会引起冲突，除非偏移量与32互质。调用内核：

int main(int argc, char* argv[]) 
{
    int offset = atoi( argv[1] ); 

    // set bank chunk to 4, just in case
    cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeFourByte);
    
    kernel<<<1, 32>>>(offset); 
    cudaDeviceSynchronize();

}

使用以下命令进行编译：nvcc bank_conflicts.cu -o bank_conflicts，我们准备好展示冲突检测了。

ncu --metrics smsp__sass_average_data_bytes_per_wavefront_mem_shared bank_conflicts 1 # ~ 99% efficiency = no conflicts :-) 
ncu --metrics smsp__sass_average_data_bytes_per_wavefront_mem_shared bank_conflicts 22 # ~ 49% efficiency = 2-way conflicts :-(
ncu --metrics smsp__sass_average_data_bytes_per_wavefront_mem_shared bank_conflicts 24 # ~ 12.5% efficiency = 8-way conflicts

作为一个额外的奖励，让我们来确立以下事实：每个银行被访问的次数为 k = GCD(offset,32) 次，因此报告的效率等于 1/k。为什么会这样呢？一个整数占据 32 位 = 4 字节，正好适合一个银行（基本切片为 4 字节）；然后线程 x 请求银行编号 bank=(x * offset)%32。这种映射使得每个值都被访问了恰好 k = GCD(offset,32) 次，这可以通过线性变换和初等数论的性质来证明 :-)

- Maciej Skorski