CUDA共享内存银行冲突报告更高

-1

我一直在努力优化一些代码，并遇到了与CUDA Nsight性能分析共享内存银行冲突报告的问题。我能够将它简化为一个非常简单的代码，Nsight报告这个代码存在银行冲突，但似乎并不存在。下面是内核：CUDA共享内存银行冲突报告更高

__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版

来源

2015-04-29 Nisrak

你的内核包含越界存储器访问。我更担心这一点，而不是银行冲突。 – talonmies

对不起，我的错误超出了界限。该数组应该是33个元素，是固定的。 – Nisrak

如果目的是按照原样运行发布的代码，并且数据类型为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，共享内存重播开销为零。

来源

2015-05-24 17:07:22

原来我的错误是数据类型为tha我正在使用。我错误地认为每个元素都会放在1个银行里。但是，双数据类型是8个字节，所以它跨越了2个共享存储体。将数据类型更改为float可以解决此问题，并且它正确显示0个银行冲突。感谢您的反馈和帮助。

来源

2015-05-01 01:35:59 Nisrak

CUDA共享内存银行冲突报告更高

回答

相关问题