GPU共享内存银行冲突

我想了解银行冲突如何发生。
如果我在全局内存中有一个大小为256的数组，并且我在一个Block中有256个线程，并且我想将该数组复制到共享内存。因此每个线程都复制一个元素。GPU共享内存银行冲突

shared_a[threadIdx.x]=global_a[threadIdx.x]

这是否会导致银行冲突？

现在该数组的大小比线程的数目较大的假设，所以现在我使用这对全局存储器复制到共享存储器：

tid = threadIdx.x; 
for(int i=0;tid+i<N;i+=blockDim.x) 
    shared_a[tid+i]=global_a[tid+i];

确实上面的代码的结果在银行冲突？

来源

2010-12-09 scatman

256是什么？字节？ – fabrizioM 2010-12-09 19:38:01

数组中的256个元素。 – scatman 2010-12-10 10:38:06

检查此问题的最佳方法是使用“Compute Visual Profiler”配置代码;这与CUDA工具包一起提供。此外还有一个很有意思的部分GPU Gems 3 - “39.2.3避免银行冲突”。

“当在同一经线访问多个线程在同一银行，发生银行冲突，除非经访问所有线程的相同的32位字在同一地址” - 第一件事，有16个记忆库，每个4字节宽。所以基本上，如果你有在半经从一个共享存储体相同的4字节读取内存的任何线程，你将有银行的冲突和系列化等

行，所以你的第一个例子：

首先让假设你的阵列是说，例如类型INT（一个32位字）的。你的代码将这些ints保存到共享内存中，在第K个线程保存到第K个内存库的任何半变形中。所以例如前半部分warp的线程0将保存到shared_a[0]这是在第一个存储体中，线程1将保存到shared_a[1]，每半个warp有16个线程映射到16个4byte bank。在下一个半翘曲中，第一个线程现在将其值保存到位于第一个存储库中的shared_a [16]中。所以如果你使用4字节这样的int，float等，那么你的第一个例子不会导致银行冲突。如果使用1字节的字（例如char），则前半部分中的经线0,1,2和3将全部将其值保存到共享内存的第一个库中，这会导致库冲突。

第二个例子：

同样，这完全取决于你所使用的字的大小，但对于例如我将使用一个4字节的字。所以寻找在第一半经：

线程数= 32

N = 64

线程0：将写入0，31，63 线程1：将写为1，32

半warp中的所有线程都并发执行，因此写入共享内存不应导致bank冲突。我必须仔细检查这一个。

希望这会有所帮助，对不起，巨大的答复！

来源

2010-12-09 09:46:04 Ljdawson

在这两种情况下，线程访问连续地址的共享内存。它取决于共享内存的元素大小，但通过线程变换连续访问共享内存不会导致“小”元素大小的银行冲突。

使用NVIDIA Visual Profiler进行剖析this code显示，对于小于32的元素大小和4的倍数（4,8,12，...，28），连续访问共享内存不会导致银行冲突。但是元素大小为32会导致银行冲突。通过Ljdawson

回答包含了一些过时的信息：

...如果你使用一个1个字节的字，如焦炭，上半年经线0，1，2和3将所有将他们的价值保存到会导致银行冲突的第一批共享内存中。

这可能是旧的GPU真正的，但对于近期的GPU与CC> = 2.x中，它们不会导致银行的冲突，有效地由于广播机制（link）。以下报价来自CUDA C PROGRAMMING GUIDE (v8.0.61) G3.3. Shared Memory。

用于经纱不会产生相同的32位字（即使两个地址落入在同一存储体）内访问的任何地址两个线程之间的存储体冲突的共享存储器请求：在这种情况下，对于读取访问，该单词被广播给请求线程（多个单词可以在单个事务中广播）并且对于写入访问，每个地址仅由其中一个线程写入（该线程执行写入是未定义的）。

这意味着，特别是，有如果字符数组如下，例如访问银行没有冲突：
extern __shared__ char shared[]; 
    char data = shared[BaseIndex + tid]; 

来源

2017-04-18 10:21:31 nglee

GPU共享内存银行冲突

回答

相关问题