0

まず、コンテキスト化を行います。CUDA で deques を使用して非ブロッキング ワーク スチール メソッドを実装しようとしています。deques (aDeques) は、グローバル メモリ内のブロック セグメント化された配列内にあり、popWork() デバイス関数には、作業をポップしてスレッドにフィードするという目的があります。グローバル deques に加えて、各ブロックには共有メモリ (aLocalStack) にスタックがあり、ローカルで動作できます。ポップは 3 レベルで発生します。最初の試行は共有スタックで、2 回目の試行はブロックが所有する両端キューで、3 回目の試行は他の両端キューのワーク スチールです。各両端キューには、グローバル メモリ配列 (aiDequesBottoms および auiDequesAges) にあるグローバル ボトム ポインターとポップ ポインターがあります。私の問題は、ブロックがグローバル deque ポインターを変更すると、GTS450 でコードをテストするときに、変更が他のブロックから見えなくなることです。キャッシュが更新されていないようです。問題が発生しない GT520 カードでもテストしました。aiDequeFlags 配列で同様の問題が発生しました。これらの問題は、揮発性を宣言することで解決されます。残念ながら、後でアトミック関数を使用する必要があるため、deque ポインター配列に対して同じことを行うことはできません。問題をより単純な例に入れずに申し訳ありませんが、この動作を再現できませんでした。この最初のスニペットには popWork() インターフェイスの説明があります。問題をより単純な例で説明できなくて申し訳ありませんが、この動作を再現できませんでした。この最初のスニペットには popWork() インターフェイスの説明があります。問題をより単純な例で説明できなくて申し訳ありませんが、この動作を再現できませんでした。この最初のスニペットには popWork() インターフェイスの説明があります。

template <int iDequeSize> //Size of each segment in aDeques 
bool __inline__ __device__ popWork(
    volatile int *aiDequeFlags , //Flags that indicates if a deque is active (has work)
    int *aiDequesBottoms , //Deque bottom pointers
    unsigned int *auiDequesAges , //Deque top pointers (29 higher bits) + 
                                  //Tag bits(3 lower bits).
    const Int2Array *aDeques , //Deques (Int2Array is an interface for 2 int arrays)
    int &uiStackBot , //Shared memory stack pointer
    int2 *aLocalStack , //Shared memory local stack
    const int &iTid , //threadIdx.x
    const int &iBid , //blockIdx.x

    //All other parameters are output

unsigned int &uiPopDequeIdx , //Choosen deque for pop
    int2 *popStartIdxAndSize , //Arrays of pop start index and sizes
    bool *bPopFlag , //Array of flags for pop in each level
unsigned int &uiActiveDequesIdx , //Flag to indicate pop failed (no more work)
    int2 &work //Actual acquired thread work)

この 2 番目のスニペットには、関数全体が含まれています。この関数を使用するカーネルは 8 ブロック、64 スレッドで起動され、最初は deque 0 だけで 1 つの作業が行われ、他のすべての deque は空です。ログを生成するための debug printf 呼び出しがいくつかあります。これは次のスニペットに表示されます。

template <int iDequeSize>
bool __inline__ __device__ popWork(volatile int *aiDequeFlags , int *aiDequesBottoms , unsigned int *auiDequesAges ,
const Int2Array *aDeques , int &uiStackBot , int2 *aLocalStack , const int &iTid , const int &iBid ,
unsigned int &uiPopDequeIdx , int2 *popStartIdxAndSize , bool *bPopFlag , unsigned int &uiActiveDequesIdx , int2 &work)
{
//Pop from local stack
if(iTid == 0)
{
    unsigned int uiAge = 0;
    bPopFlag[0] = popBottom(uiStackBot , uiAge , popStartIdxAndSize[iBid]); 
    bPopFlag[3] = bPopFlag[0];
}

__syncthreads();

if(bPopFlag[0])
{
    if(iTid < popStartIdxAndSize[iBid].y)
    {
        work = aLocalStack[popStartIdxAndSize[iBid].x + iTid];
    }
}
else
{
    if(iTid == 0)
    {   //Try to pop from block deque

        bPopFlag[1] = popBottom(aiDequesBottoms[iBid] , auiDequesAges[iBid] , popStartIdxAndSize[iBid]);

        if(bPopFlag[1])
        {
            uiPopDequeIdx = iBid;
            //Debug
            if(iBid == 0)
            {
                printf("Block %d pop global deque. Bottom=%d\n" , iBid , aiDequesBottoms[iBid]);
            }
            //
        }
        else
        {
            aiDequeFlags[iBid] = 0;
            popStartIdxAndSize[iBid].x = INFTY;
            uiPopDequeIdx = INFTY;
        }
        bPopFlag[3] = bPopFlag[1];
        bPopFlag[2] = false;
    }
    __syncthreads();

    if(!bPopFlag[1])
    {
        //Verify if lazy steal can be done.
        if(iTid < NDEQUES)
        {
            if(popStartIdxAndSize[iTid].x != INFTY && iTid != iBid)
            {
                atomicMin(&uiPopDequeIdx , iTid);
                bPopFlag[2] = true;
                bPopFlag[3] = true;
            }
        }

        __syncthreads();

        if(iTid == uiPopDequeIdx)
        {
            popStartIdxAndSize[iBid] = popStartIdxAndSize[iTid];
            popStartIdxAndSize[iTid].x = INFTY;
        }

        while(!bPopFlag[3])
        {   //No more work, try to steal some!
            __syncthreads();

            if(iTid == 0)
            {
                uiActiveDequesIdx = 0;
            }
            __syncthreads();

            if(iTid < NDEQUES)
            {
                if(aiDequeFlags[iTid] == 1)
                {
                    uiActiveDequesIdx = 1;

                    //Debug
                    printf("Block %d steal attempt on block %d. Victim bottom=%d\n" , blockIdx.x , threadIdx.x , aiDequesBottoms[iTid]);
                    //

                    if(popTop(aiDequesBottoms , auiDequesAges , iTid , popStartIdxAndSize[iTid]))
                    {
                        aiDequeFlags[iBid] = 1;
                        atomicMin(&uiPopDequeIdx , iTid);
                        bPopFlag[3] = true;

                        //Debug
                        //printf("%d ss %d %d %d\n" , iBid , iTid , popStartIdxAndSize[iTid].x , popStartIdxAndSize[iTid].y);
                        //
                    }
                }
            }

            __syncthreads();

            if(uiActiveDequesIdx == 0)
            { //No more work to steal. End.
                break;
            }

            if(iTid == uiPopDequeIdx)
            {
                popStartIdxAndSize[iBid] = popStartIdxAndSize[iTid];
                popStartIdxAndSize[iTid].x = INFTY;
            }

            __syncthreads();
        }
    }

    __syncthreads();

    if(bPopFlag[3] && iTid < popStartIdxAndSize[iBid].y) //assuming number of threads >= WORK_SIZE
    {
        aDeques->getElement(work , uiPopDequeIdx*iDequeSize + popStartIdxAndSize[iBid].x + iTid);
    }
}

return bPopFlag[3];

}

この最後のスニペットは、生成されたログです。プッシュ ライン ("Block X push. Bottom=Y") は、ここには示されていなかったプッシュ関数によって生成されました。最初は、ブロック 0 だけに 1 つの作業があることに注意してください。

Block 0 pop global deque. Bottom=0
Block 4 steal attempt on block 0. Victim bottom=0
Block 2 steal attempt on block 0. Victim bottom=1
Block 6 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 0. Victim bottom=1
Block 3 steal attempt on block 0. Victim bottom=1
Block 7 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 0. Victim bottom=1
Block 6 steal attempt on block 0. Victim bottom=1
Block 3 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 0. Victim bottom=1
Block 4 steal attempt on block 0. Victim bottom=0
Block 7 steal attempt on block 0. Victim bottom=1
Block 0 push. Bottom=448
Block 2 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 0. Victim bottom=1
Block 6 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 0. Victim bottom=1
Block 3 steal attempt on block 0. Victim bottom=1
Block 7 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 0. Victim bottom=1
Block 2 steal attempt on block 4. Victim bottom=0
Block 1 steal attempt on block 0. Victim bottom=1
Block 1 steal attempt on block 4. Victim bottom=0
Block 5 steal attempt on block 0. Victim bottom=1
Block 5 steal attempt on block 4. Victim bottom=0
Block 4 push. Bottom=384

ご覧のとおり、ブロック 4 のみがブロック 0 の両端キューのボトム ポインターの変更を確認できます。ポインターの変更後に __threadfence() 呼び出しをいくつか追加しようとしましたが、成功しませんでした。注目してくれてありがとう！

4

2 に答える 2

3

唯一の有効な解決策は L1 キャッシュをオフにすることであるというコメントに基づいているようです。これは、コンパイル時に次のスイッチを nvcc に渡すことにより、プログラム全体で実現できます。

–Xptxas –dlcm=cg

L1 キャッシュは、デバイス全体ではなく、SM のプロパティ/リソースです。スレッドブロックは特定の SM で実行されるため、L1 キャッシュ内の 1 つのスレッドブロックのアクティビティは、別のスレッドブロックおよびその L1 キャッシュのアクティビティと一貫性がない可能性があります (別の SM で実行されていると仮定します)。グローバルメモリ内の場所。異なる SM の L1 キャッシュは相互に接続されておらず、相互の一貫性が保証されていません。

L2 キャッシュはデバイス全体に適用されるため、個々のスレッドブロックの観点からは「一貫性」があることに注意してください。L1 キャッシングをオフにしても L2 キャッシングには影響がないため、キャッシングの利点が得られる可能性はまだありますが、L2 からのリクエストを満たすために必要な時間は、L1 からのリクエストを満たすために必要な時間よりも長いため、オフにします。プログラム全体の L1 キャッシングは、物事を機能させるための非常に大きなハンマーです。

変数定義のvolatile前にあるキーワードには、ロード時に L1 キャッシュをスキップするようにコンパイラに指示する効果があるはずです (私の理解によると)。しかし、volatile 自体は書き込みパスをアドレス指定しないため、1 つの SM 内の 1 つのスレッドブロックがvolatile読み取りを実行し、L2 から値を引き出し、その値を変更してから書き戻すことが可能であり、最終的に L1 に到達します (追い出されるまで）。別のスレッドブロックが同じグローバル値を読み取る場合、更新の効果が見られない場合があります。

面倒な__threadfence()を入念に使用すると、そのような更新を L1 から L2 に強制的に実行して、他のスレッドブロックがそれらを読み取れるようにする必要があります。ただし、これにより、値が書き込まれたときから、他の SM/スレッドブロックによって監視可能になるまでの同期ギャップが残ります。

(グローバル)アトミックには、使用される値を読み書きするために「グローバル メモリ」に直接移動する効果も必要です。

コードを調べて、グローバルに同期された場所からのすべての可能な読み取りが適切に処理されること (たとえば、volatileアトミックを使用またはアトミックを使用) と、グローバルに同期された場所へのすべての可能な書き込みが適切に処理されること (たとえば、__threadfence()またはアトミックを使用) を確認することも有益な場合があります。 、異なるブロック間の競合状態もチェックします。

発見されたように、GPU 内でグローバルに同期された安定した環境を作成するプロセスは自明ではありません。これらの他の質問も興味深いかもしれません (例えばKepler に関して) (そして例えばグローバルセマフォの議論)。

編集: コメントに投稿された質問に答えるには、次のように言います。

多分、問題ないです。ただし__threadfence()、最大完了時間については (私が知っている) 保証はありません。したがって、グローバル ロケーションが更新された時点で、実行中のスレッドブロック/SM に関連付けられている L1 のみが更新されます。次に、 をヒットし__threadfence()ます。おそらく、スレッドフェンスが完了するまでに時間がかかり、この間に別のスレッドブロックが同じ SM に常駐し、実行のために持ち込まれ (前のスレッド/ワープ/ブロックがスレッドフェンスで停止している間)、更新されたグローバル値を「見る」ことができます。その SM に関連付けられた (ローカル) L1 内。他の SM で実行されている他のスレッドブロックには、「古い」値が表示されます。__threadfence()完了します。これは、私が「同期ギャップ」の可能性と呼んでいるものです。2 つの異なるブロックは、短時間の間、2 つの異なる値を見ることができます。これが問題になるかどうかは、グローバル値がブロック間の同期にどのように使用されているかに依存します (これが議論中のトピックであるため)。および同期用の書き込みパス。

編集＃2：コメントから、アトミックプラスの使用の組み合わせvolatileも問題を解決したようです。

于 2013-01-25T15:00:32.247 に答える
0

率直に言って、あなたのコードはインデックスで過度に複雑であり、さらに重要なことに、不完全です。どのようpopBottomにpopTop機能しますか？さらに、push操作はどのように実装されますか？これら 2 つは、正しく動作し、同期の問題が発生しないようにするために、慎重に作成する必要があります。

たとえば、あるブロックがグローバル メモリ キューに何かをプッシュしようとしているときに、別のブロックが同じ瞬間にそこから読み取ろうとするとどうなるでしょうか? これは非常に重要であり、正しく行われないと、非常にまれな状況でクラッシュする可能性があります。たとえば、まだ書き込まれていないデータ セルからポップする場合があります。

同様のことを実装していたとき、つまりすべてのブロック間で共有される単一のグローバル メモリ デュークを実装していたとき、さらに各データ セルを空、占有、無効としてマークしていました。疑似コードでは、アルゴリズムは多かれ少なかれ次のように機能しました。

/* Objects of this class should reside in CUDA global memory */
template <typename T, size_t size>
class WorkQueue {
private:
    size_t head, tail;
    size_t status[size];
    T data[size];

    enum {
        FieldFree = 0,
        FieldDead = 1,
        FieldTaken = 2
    };      

public:
    /* 
       This construction should actually be done by host on the device,
       before the actual kernel using it is launched!
       Zeroing the memory should suffice.
    */
    WorkQueue() : head(0), tail(0) {
        for (size_t i=0; i<size; ++i)
            status[i]=FieldFree;
    }   

    __device__ bool isEmpty() { return head==tail; }

    /* single thread of a block should call this */
    __device__ bool push(const T& val) {
        size_t oldFieldStatus;
        do {
            size_t cell = atomicInc(&tail,size-1);
            data[cell]=val;
            __threadfence(); //wait untill all blocks see the above change
            oldFieldStatus=atomicCAS(&status[cell],FieldFree,FieldTaken); //mark the cell as occupied
        } while (oldFieldStatus!=FieldFree); 
        return true;
    }

    /* single thread of a block should call this */
    __device__ bool pop(T& out) {
        size_t cellStatus;
        size_t cell;
        do {
            cell=atomicInc(&head,size-1);
            cellStatus=atomicCAS(&status[cell],FieldFree,FieldDead);
            //If cell was free, make it dead - any data stored there will not be processed. Ever.
        } while (cellStatus==FieldDead);
        if (cellStatus!=FieldTaken)
            return false;
        out = data[cell];
        status[cell]=FieldFree;
        return true;
    }
};

セルのステータスなしで実装する信頼できる方法がわかりません。そうしないと、2 つの異なるブロックの 2 つのスレッドがデキューの同じセルにプッシュ/ポップしようとすると、悪いことが起こります。上記のアプローチでは、ポップ スレッドがポップに失敗し、false を返し、セルを としてマークしdead、プッシュ スレッドが次のセルへのプッシュを再試行するという最悪のケースが発生する可能性があります。背後にある考え方は、ポップ スレッドがポップに失敗した場合、とにかく行うべき作業があまりなく、ブロックが終了する可能性があるということです。そのアプローチでは、並行して実行されているブロックと同じ数のセルのみを「殺す」ことができます。

上記のコードでは、オーバーフローをチェックしていないことに注意してください。

于 2013-01-23T18:54:35.543 に答える