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

【问题标题】：CUDA shared memory bank conflicts report higherCUDA 共享内存库冲突报告更高
【发布时间】：2015-04-29 03:47:36
【问题描述】：

我一直致力于优化一些代码，但在使用 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 版

【问题讨论】：

您的内核包含越界内存访问。比起银行冲突，我更担心这一点。
对不起，我的错误越界了。数组应该是 33 个元素，固定的。

标签： cuda gpu shared-memory bank-conflict

【解决方案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，共享内存重播开销为零。

【讨论】：

【解决方案2】：

原来我的错误在于我使用的数据类型。我错误地认为每个元素都将放在 1 个银行中是理所当然的。但是，双精度数据类型是 8 字节，因此它跨越 2 个共享内存库。将数据类型更改为浮点数解决了这个问题，它正确显示了 0 个银行冲突。感谢您的反馈和帮助。

【讨论】：