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共享內存銀行衝突報告更高

回答

相關問題