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
• 库达7
• Nsight Visual Studio版本4.5

Answer 1

如果目的是使用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规范，共享内存重放开销为零。

Answer 2

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