-1

我一直在优化一些代码，并在 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 时遇到过问题，或者我只是在银行冲突的运作方面遗漏了一些东西？我将不胜感激任何反馈！

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

  • 视窗 8
  • GTX 770
  • Visual Studio 社区 2013
  • CUDA 7
  • Nsight Visual Studio 4.5版
4

2 回答 2

3

如果意图是按原样运行发布的代码，使用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-24T17:07:22.363 回答
0

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

于 2015-05-01T01:35:59.413 回答