2

グローバルな書き込みが表示されることを保証するのは私の理解です (たとえば、ポインタを揮発性として宣言せずに CUDA グローバル メモリ コヒーレンスを適用するにはどうすればよいですか?、 GTS 250 デバイスとFermi デバイス間の CUDA ブロック同期の違い、およびnvidia 開発者ゾーンのこの投稿を参照してください)。__threadfence()スレッドが続行される前に、他のスレッドに。ただし、別のスレッドは、が返された後でも、その L1 キャッシュから古い値を読み取る可能性があります。__threadfence()

あれは：

スレッド A は、一部のデータをグローバル メモリに書き込み、次に を呼び出します__threadfence()。次に、 が戻った後、書き込みが他の すべて__threadfence()のスレッドから見えるようになると、スレッド B はこのメモリ位置から読み取るように要求されます。L1 にデータがあることがわかったので、それをロードします。開発者にとって残念なことに、スレッド B の L1 のデータは古くなっています (つまり、スレッド A がこのデータを更新する前と同じです)。

まず第一に、これは正しいですか？

そうだとすると、データがL1にないことが確実な場合（ややありそうにない？）、または読み取りが常にL1をバイパスする場合（揮発性またはアトミックなど）__threadfence()にのみ有用であるように思えます。これは正しいです？

私は比較的単純なユースケースを持っているので質問します-二分木にデータを伝播する-アトミックに設定されたフラグとを使用し__threadfence()ます:ノードに到達する最初のスレッドが終了し、2番目のスレッドが2つの子に基づいてそれにデータを書き込みます(たとえば、最小限のデータ)。これはほとんどのノードで機能しますが、通常、少なくとも 1 つのノードで失敗します。データを宣言するとvolatile一貫して正しい結果が得られますが、L1 から古い値が取得されない 99% 以上のケースでパフォーマンス ヒットが発生します。これがこのアルゴリズムの唯一の解決策であることを確認したいと思います。簡単な例を以下に示します。ノード配列は幅優先で並べられていることに注意してください。リーフは index から始まりstart、既にデータが入力されています。

__global__ void propagate_data(volatile Node *nodes,
                               const unsigned int n_nodes,
                               const unsigned int start,
                               unsigned int* flags)
{
    int tid, index, left, right;
    float data;
    bool first_arrival;

    tid = start + threadIdx.x + blockIdx.x*blockDim.x;

    while (tid < n_nodes)
    {
        // We start at a node with a full data section; modify its flag
        // accordingly.
        flags[tid] = 2;

        // Immediately move up the tree.
        index = nodes[tid].parent;
        first_arrival = (atomicAdd(&flags[index], 1) == 0);

        // If we are the second thread to reach this node then process it.
        while (!first_arrival)
        {
            left = nodes[index].left;
            right = nodes[index].right;

            // If Node* nodes is not declared volatile, this occasionally
            // reads a stale value from L1.
            data = min(nodes[left].data, nodes[right].data);

            nodes[index].data = data;

            if (index == 0) {
                // Root node processed, so all nodes processed.
                return;
            }

            // Ensure above global write is visible to all device threads
            // before setting flag for the parent.
            __threadfence();

            index = nodes[index].parent;
            first_arrival = (atomicAdd(&flags[index], 1) == 0);
        }
        tid += blockDim.x*gridDim.x;
    }
    return;
}
4

1 に答える 1