0

Nvidia の開発者サイトで OpenCL カーネル サンプル コードの一部を見つけました。目的の関数maxOneBlockは、配列の最大値を見つけてmaxValuemaxValue[0] に格納することです。

ループ部分については完全に理解していましたが、unroll次の部分については混乱していました。各ステップが完了した後にアンロール部分でスレッドを同期する必要がないのはなぜですか?

例: 1 つのスレッドが localId と localId+32 の比較を行った場合、他のスレッドがその結果を localId+16 に保存したことをどのように確認しますか?

カーネルコード:

void maxOneBlock(__local float maxValue[],
                 __local int   maxInd[])
{
    uint localId   = get_local_id(0);
    uint localSize = get_local_size(0);
    int idx;
    float m1, m2, m3;

    for (uint s = localSize/2; s > 32; s >>= 1)
    {
        if (localId < s) 
        {
            m1 = maxValue[localId];
            m2 = maxValue[localId+s];
            m3 = (m1 >= m2) ? m1 : m2;
            idx = (m1 >= m2) ? localId : localId + s;
            maxValue[localId] = m3;
            maxInd[localId] = maxInd[idx];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // unroll the final warp to reduce loop and sync overheads
    if (localId < 32)
    {
        m1 = maxValue[localId];
        m2 = maxValue[localId+32];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 32;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];


        m1 = maxValue[localId];
        m2 = maxValue[localId+16];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 16;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+8];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 8;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+4];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 4;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+2];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 2;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];

        m1 = maxValue[localId];
        m2 = maxValue[localId+1];
        m3 = (m1 > m2) ? m1 : m2;
        idx = (m1 > m2) ? localId : localId + 1;
        maxValue[localId] = m3;
        maxInd[localId] = maxInd[idx];
    }
}
4

1 に答える 1

0

各ステップが完了した後にアンロール部分がスレッドを同期する必要がないのはなぜですか?

サンプルが正しくありません。実際には、各ステップの後にバリアが必要です。

サンプルは、NVIDIA ハードウェアのワープの基本的な実行メカニズムを利用する方法であるワープ同期スタイルで記述されているように見えますが、基本的な実行メカニズムが変更された場合、またはコンパイラの最適化が存在する場合、不適切な同期によりサンプルが中断されます。 .

于 2015-06-01T05:04:18.803 に答える