我正在尝试理解银行冲突是如何发生的。
我在全局内存中有一个大小为256的数组，并且我有一个单一块中的256个线程，我想将该数组复制到共享内存。因此，每个线程都复制一个元素。
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];
上述代码会导致银行冲突吗？
最好的检查方法是使用“Compute Visual Profiler”对代码进行分析，这是CUDA Toolkit附带的工具。此外，在GPU Gems 3中有一个很好的章节，“39.2.3避免Bank冲突”。
“当同一warp中的多个线程访问同一bank时，除非warp的所有线程都在同一32位字内访问同一地址，否则会发生bank冲突”-首先，共有16个内存bank，每个bank宽度为4个字节。因此，如果半个warp中的任何一个线程从共享内存bank中读取相同4个字节的内存，就会发生bank冲突并出现序列化等问题。
好的，现在看看第一个例子：
首先假设你的数组是例如int（一个32位字）类型的。你的代码将这些int保存到共享内存中，每个半个warp的Kth线程都会保存到Kth内存bank。例如，第一个半个warp的第0个线程将保存到shared_a[0]中，它位于第一个内存bank中，线程1将保存到shared_a[1]中，每个半个warp有16个线程，它们映射到16个4字节的bank上。在下一个半个warp中，第一个线程现在将其值保存到shared_a [16]中，这又是在第一个内存bank中。因此，如果您使用一个4字节的word（如int，float等），那么您的第一个例子将不会导致bank冲突。但是如果您使用一个1字节的word（例如char），则在第一个半个warp中，线程0、1、2和3都会将它们的值保存到共享内存的第一个bank中，这将导致bank冲突。
第二个例子：
同样，这将取决于您使用的字的大小，但为了举例，我将使用4字节的word。所以看一下第一个半个warp：
线程数 = 32
N = 64
线程0：将写入0、31、63；线程1：将写入1、32。
所有半warp中的线程并发执行，因此对共享内存的写操作不应该导致bank冲突。我需要再仔细检查一下这个问题。
希望这可以帮到您，非常抱歉回复有些长！
Ljdawson的回答包含一些过时的信息：
... 如果您使用1字节的单词（例如char），在第一半warp线程0、1、2和3将所有值保存到共享内存的第一个bank中，这将导致bank冲突。
对于具有cc >= 2.x的最新GPU来说，这可能对旧GPU是正确的，但它们不会引起bank冲突，有效地由于广播机制(link)。以下引用来自CUDA C编程指南 (v8.0.61) G3.3. Shared Memory。
A shared memory request for a warp does not generate a bank conflict between two threads that access any address within the same 32-bit word (even though the two addresses fall in the same bank): In that case, for read accesses, the word is broadcast to the requesting threads (multiple words can be broadcast in a single transaction) and for write accesses, each address is written by only one of the threads (which thread performs the write is undefined).
This means, in particular, that there are no bank conflicts if an array of char is accessed as follows, for example:
extern __shared__ char shared[]; char data = shared[BaseIndex + tid];
for (int i = tid ; i < N ; i += blockDim.x) shared_a[i] = global_a[i];。 - Tom