如何减少此代码中的银行冲突?
Posted
技术标签:
【中文标题】如何减少此代码中的银行冲突?【英文标题】:How can I diminish bank conflicts in this code? 【发布时间】:2012-12-18 01:16:11 【问题描述】:当 Nsight 分析时,这段 CUDA 代码报告了很多银行冲突。第一个 sn-p 包含常量定义和内核调用:
// Front update related constants
#define NDEQUES 6
#define FRONT_UPDATE_THREADS 480
#define BVTT_DEQUE_SIZE 500000
#define FRONT_DEQUE_SIZE 5000000
#define FRONT_UPDATE_SHARED_SIZE FRONT_UPDATE_THREADS*2
updateFront<OBBNode , OBB , BVTT_DEQUE_SIZE , FRONT_DEQUE_SIZE , FRONT_UPDATE_THREADS>
<<<NDEQUES, FRONT_UPDATE_THREADS>>>
(d_aFront , d_aOutputFront , d_aiFrontCounts , d_aWorkQueues , d_aiWorkQueueCounts , d_collisionPairs ,
d_collisionPairIndex , obbTree1 , d_triIndices1);
第二个sn-p有内核代码:
template<typename TreeNode , typename BV , unsigned int uiGlobalWorkQueueCapacity , unsigned int uiGlobalFrontCapacity ,
unsigned int uiNThreads>
void __global__ updateFront(Int2Array *aFront , Int2Array *aOutputFront , int *aiFrontIdx , Int2Array *aWork_queues ,
int* aiWork_queue_counts , int2 *auiCollisionPairs , unsigned int *uiCollisionPairsIdx , const TreeNode* tree ,
uint3 *aTriIndices)
__shared__ unsigned int uiInputFrontIdx;
__shared__ unsigned int uiOutputFrontIdx;
__shared__ unsigned int uiWorkQueueIdx;
__shared__ int iLeafLeafOffset;
__shared__ int iNode0GreaterOffset;
__shared__ int iNode1GreaterOffset;
__shared__ int aiLeafLeafFrontX[FRONT_UPDATE_SHARED_SIZE];
__shared__ int aiLeafLeafFrontY[FRONT_UPDATE_SHARED_SIZE];
__shared__ int aiNode0GreaterFrontX[FRONT_UPDATE_SHARED_SIZE];
__shared__ int aiNode0GreaterFrontY[FRONT_UPDATE_SHARED_SIZE];
__shared__ int aiNode1GreaterFrontX[FRONT_UPDATE_SHARED_SIZE];
__shared__ int aiNode1GreaterFrontY[FRONT_UPDATE_SHARED_SIZE];
if(threadIdx.x == 0)
uiInputFrontIdx = aiFrontIdx[blockIdx.x];
uiOutputFrontIdx = 0;
uiWorkQueueIdx = aiWork_queue_counts[blockIdx.x];
iLeafLeafOffset = 0;
iNode0GreaterOffset = 0;
iNode1GreaterOffset = 0;
__syncthreads();
unsigned int uiThreadOffset = threadIdx.x;
while(uiThreadOffset < uiInputFrontIdx + FRONT_UPDATE_THREADS - (uiInputFrontIdx % FRONT_UPDATE_THREADS))
if(uiThreadOffset < uiInputFrontIdx)
int2 bvttNode;
aFront->getElement(bvttNode , blockIdx.x*FRONT_DEQUE_SIZE + uiThreadOffset);
TreeNode node0 = tree[bvttNode.x];
TreeNode node1 = tree[bvttNode.y];
if(node0.isLeaf() && node1.isLeaf())
int iOffset = atomicAdd(&iLeafLeafOffset , 1);
//Bank conflict source
aiLeafLeafFrontX[iOffset] = bvttNode.x;
aiLeafLeafFrontY[iOffset] = bvttNode.y;
//End of bank conflict source
else if(node1.isLeaf() || (!node0.isLeaf() && (node0.bbox.getSize() > node1.bbox.getSize())))
// node0 is bigger. Subdivide it.
int iOffset = atomicAdd(&iNode0GreaterOffset , 1);
//Bank conflict source
aiNode0GreaterFrontX[iOffset] = bvttNode.x;
aiNode0GreaterFrontY[iOffset] = bvttNode.y;
//End of bank conflict source
else
// node1 is bigger. Subdivide it.
int iOffset = atomicAdd(&iNode1GreaterOffset , 1);
//Bank conflict source
aiNode1GreaterFrontX[iOffset] = bvttNode.x;
aiNode1GreaterFrontY[iOffset] = bvttNode.y;
//End of bank conflict source
__syncthreads();
/* ... */
uiThreadOffset += uiNThreads;
__syncthreads();
我想知道为什么会发生银行冲突。我认为可能发生冲突的唯一方法是,如果映射到同一银行的不同数组中的访问被序列化。
【问题讨论】:
在不知道共享内存数组的维度的情况下,很难评论银行冲突...... 我已经编辑了内核调用代码和常量定义。 【参考方案1】:我看到了两种可能性。需要进一步测试才能选择哪一个是罪魁祸首:
银行冲突不是发生在您选择的位置,而是发生在atomicAdd
操作上,这些操作也适用于共享内存。我相信 shmem 上的原子也可以增加内部冲突计数器。 (信念未经检验!)
您遇到了两个或多个 warp 以原子方式增加相同值的情况 - 这可能在同时运行 2 或 4 个 warp 的较新硬件上发生。 (也需要测试来确认或否认这一点)。因此,一个warp 中的线程实际上可能会获得相当遥远的iOffset
值,并且您最终会遇到一些随机的银行冲突。
但是,如果以上任何一个都属实,我就不会太担心这些冲突。在第一种情况下 - atomicAdd
无论如何都会影响你的表现。在后一种情况下,我不希望经常发生超过 2 路银行冲突。除非你遇到一些非常罕见的极端情况......
【讨论】:
我已经做了一个测试,看看冲突是否因为原子而发生。我已经评论了所有共享数组访问,因此只有原子操作访问了共享内存。在这种情况下,Nsight 报告了 0 次银行冲突。 你是怎么做到的?通过将 aiLeafLeafFrontX(和其他)数组移动到全局?还是直接扔掉?请记住,此类更改可能会影响 warp 调度并减少其他地方的银行冲突... 我刚刚注释了访问共享数组的代码,只留下了计算数组偏移量并运行 Nsight CUDA 分析的原子。例如:' int iOffset = atomicAdd(&iLeafLeafOffset , 1); //aiLeafLeafFrontX[iOffset] = bvttNode.x; //aiLeafLeafFrontY[iOffset] = bvttNode.y; ' 我不明白为什么 warp 调度会影响银行冲突。我的理解是冲突与不依赖于调度的数组索引有关。能详细解释一下吗? Warp 调度不会直接影响 bank 冲突,但可能会影响 atomics 值的分布。此外 - 如果您只是注释掉一部分代码,编译器可能会检测到剩余的代码是多余的,并删除您认为的更多代码。特别是 - 如果您从不写入全局内存,您的内核可能会缩减为空内核。 你是对的。调度会影响线程执行原子操作的顺序,并可能产生存储库冲突。我将尝试使用前缀和方法来更好地控制索引。我会将您的答案标记为正确。感谢您的提示!以上是关于如何减少此代码中的银行冲突?的主要内容,如果未能解决你的问题,请参考以下文章