0

この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);

2番目のスニペットにはカーネルコードが含まれています。

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

なぜ銀行の対立が起こっているのか知りたい。競合が発生する可能性があると思う唯一の方法は、同じバンクにマップする異なるアレイのアクセスがシリアル化された場合です。

4

1 に答える 1

1

2つの可能性があります。どちらが原因であるかを選択するには、さらにテストが必要です。

  • バンクの競合は、選択した場所からではなく、atomicAdd共有メモリでも機能する操作から発生しています。shmemのアトミックは、内部の競合カウンターも増やすことができると思います。（信念はテストされていません！）

  • 2つ以上のワープが同じ値をアトミックに増加させている状況に遭遇しました。これは、2つまたは4つのワープを同時に実行する新しいハードウェアで発生する可能性があります。（これも確認または拒否するためにテストが必要です）。その結果、1つのワープ内のスレッドは実際にはかなり離れたiOffset値を取得する可能性があり、ランダムなバンク競合が発生することになります。

ただし、上記のいずれかが当てはまる場合は、競合についてあまり心配する必要はありません。最初のケースでは-atomicAddとにかくあなたのパフォーマンスに影響を与えます。後者の場合、双方向の銀行の競合が頻繁に発生することはないと思います。本当に珍しいコーナーケースにぶつからない限り…。

于 2013-01-03T21:35:31.880 に答える