CUDA 共享内存库冲突报告更高

Question

我一直致力于优化一些代码，运行解决了与 CUDA Nsight 性能分析共享内存库冲突报告的问题。我能够将它简化为一段非常简单的代码，Nsight 将其报告为存在银行冲突，而实际上它似乎不应该存在。下面是内核：

__global__ void conflict() {
    __shared__ double values[33];
    values[threadIdx.x] = threadIdx.x;
    values[threadIdx.x+1] = threadIdx.x;
}

调用它的主要函数：

int main() {
    conflict<<<1,32>>>();
}

请注意，我正在使用单个扭曲来真正将其减少到最低限度。当我运行代码时，Nsight 说有 1 个银行冲突，但根据我阅读的所有内容，应该没有。对于共享内存数组的每次访问，每个线程都在访问连续的值，每个值属于不同的存储区。

有没有其他人遇到过 Nsight 报告的问题，或者我只是遗漏了一些关于银行冲突功能的问题？如果有任何反馈，我将不胜感激！

顺便说一句，我正在运行进行以下设置：

Windows 8
GTX 770
Visual Studio 社区 2013
CUDA 7
Nsight Visual studio 版版本 4.5

Answer 1

事实证明我的错误在于我使用的数据类型。我错误地认为每个元素都将放在 1 个库中。但是，double 数据类型是 8 个字节，因此它跨越 2 个共享内存组。将数据类型更改为 float 解决了这个问题，它正确地显示了 0 个银行冲突。感谢您的反馈和帮助。

Answer 2

如果目的是运行按原样发布代码，使用 double 数据类型，并且没有银行冲突，我相信适当使用 cudaDeviceSetSharedMemConfig 是可能的（在 cc3.x 台设备上）。这是一个测试用例：

$ cat t750.cu
#include <stdio.h>

typedef double mytype;


template <typename T>
__global__ void conflict() {
    __shared__ T values[33];
    values[threadIdx.x] = threadIdx.x;
    values[threadIdx.x+1] = threadIdx.x;
}

int main(){

#ifdef EBM
  cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
#endif

  conflict<mytype><<<1,32>>>();
  cudaDeviceSynchronize();
}

$ nvcc -arch=sm_35 -o t750 t750.cu
t750.cu(8): warning: variable "values" was set but never used
          detected during instantiation of "void conflict<T>() [with T=mytype]"
(19): here

$ nvprof --metrics shared_replay_overhead ./t750
==46560== NVPROF is profiling process 46560, command: ./t750
==46560== Profiling application: ./t750
==46560== Profiling result:
==46560== Metric result:
Invocations                               Metric Name                        Metric Description         Min         Max         Avg
Device "Tesla K40c (0)"
 Kernel: void conflict<double>(void)
          1                    shared_replay_overhead             Shared Memory Replay Overhead    0.142857    0.142857    0.142857
$ nvcc -arch=sm_35 -DEBM -o t750 t750.cu
t750.cu(8): warning: variable "values" was set but never used
          detected during instantiation of "void conflict<T>() [with T=mytype]"
(19): here

$ nvprof --metrics shared_replay_overhead ./t750
==46609== NVPROF is profiling process 46609, command: ./t750
==46609== Profiling application: ./t750
==46609== Profiling result:
==46609== Metric result:
Invocations                               Metric Name                        Metric Description         Min         Max         Avg
Device "Tesla K40c (0)"
 Kernel: void conflict<double>(void)
          1                    shared_replay_overhead             Shared Memory Replay Overhead    0.000000    0.000000    0.000000
$

使用 EightByteMode 的规范，共享内存重放开销为零。

CUDA 共享内存库冲突报告更高

CUDA shared memory bank conflicts report higher

cuda

gpu

shared-memory

bank-conflict