首页 文章

2.x设备中的银行冲突

提问于
浏览 次
1

具有2.x设备的设备中的银行冲突是什么？据我了解CUDA C编程指南，在2.x设备中，如果两个线程在同一共享内存库中访问相同的32位字，则不会导致存储体冲突 . 相反，这个词是广播的 . 当两个线程在同一共享内存库中写入相同的32位字时，只有一个线程成功 .

由于片上存储器为64 KB（共享存储器为48 KB，L1为16 KB，反之亦然），并且它由32个存储体组成，我假设每个存储体由2 KB组成 . 所以我认为如果两个线程在同一个共享内存库中访问两个不同的32位字，就会出现银行冲突 . 它是否正确？

cuda gpu shared-memory bank-conflict

1 回答

  • 3

    你的描述是正确的 . 有许多访问模式可以产生银行冲突，但这里有一个简单而常见的例子：跨步访问 .

    __shared__ int smem[512];

int tid = threadIdx.x;

x = smem[tid * 2]; // 2-way bank conflicts
y = smem[tid * 4]; // 4-way bank conflicts
z = smem[tid * 8]; // 8-way bank conflicts
// etc.

    Bank ID = index％32，所以如果你看一下x，y和z访问中的地址模式，你可以看到在32个线程的每个warp中，对于x，2个线程将访问每个bank，对于y， 4个线程将访问每个银行，对于z，8个线程将访问每个银行 .

    回复于

相关问题