如何減少此代碼中的銀行衝突？

Question

這段CUDA代碼在Nsight分析時會報告很多銀行衝突。第一個片段包含常量的定義和內核調用：

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

我想知道爲什麼銀行衝突ocurring。我認爲衝突可能發生的唯一方法是如果映射到同一個銀行的不同陣列中的訪問被序列化。

Answer 1

我看到兩種可能性。進一步的測試需要選擇哪一個是culpit：

• 銀行衝突不是從你選擇的地點發生的，但是從atomicAdd操作也共享內存工作。我相信shmem上的原子可以增加內部衝突計數器。 （這個信念沒有經過測試！）

• 你遇到兩個或兩個以上的經線原子上增加相同的值的情況 - 這可能是在同時運行2或4個經線的新硬件上的可能性。 （測試也需要確認或否認）。因此，一個warp內的線程實際上可能會得到相當遠的值，並且最終會產生一些隨機的bank衝突。

但是，如果有一個以上是真的，我woulndn't擔心衝突了。在第一種情況下 - 無論如何，atomicAdd都會打你的表演。在後一種情況下，我不會指望經常會出現雙向銀行衝突。除非你打一些非常罕見的極端情況....