如何在代码中减少银行冲突?

Question

这段CUDA代码在被insight分析时报告了许多银行冲突。第一个代码片段包含常量定义和内核调用:

// 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);

第二个代码片段包含内核代码:

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();
}

我想知道为什么银行冲突会发生。我认为唯一可能发生冲突的方式是，如果映射到同一银行的不同数组中的访问被序列化。

Answer 1

我认为有两种可能性。需要进一步的测试来选择哪一个是罪魁祸首:

• 银行冲突不是发生在您选择的位置，而是发生在也在共享内存上工作的atomicAdd操作。我相信shmem上的原子也会增加内部冲突计数器。(信念未经检验!)

• 你遇到了一个情况，两个或更多的翘曲自动增加相同的值-这可能是在较新的硬件上运行2或4次翘曲同时发生的可能性。(也需要测试来证实或否认这一点)。因此，在一个经线内的线程实际上可能会得到相当远的iOffset值，你最终会有一些随机的银行冲突。

然而，如果以上任何一个是真的，我不会太担心冲突。在第一种情况下- atomicAdd无论如何都会影响您的性能。在后一种情况下，我不认为经常会有超过双向的银行冲突。除非遇到罕见的情况....