2

CUDA アトミック API に atomicLoad 関数がないという問題に直面しました。stackoverflow を検索した後、次の CUDA atomicLoadの実装を見つけました

しかし、次の例では、この関数が機能していないようです。

#include <cassert>
#include <iostream>
#include <cuda_runtime_api.h>

template <typename T>
__device__ T atomicLoad(const T* addr) {
    const volatile T* vaddr = addr;  // To bypass cache
    __threadfence();                 // for seq_cst loads. Remove for acquire semantics.
    const T value = *vaddr;
    // fence to ensure that dependent reads are correctly ordered
    __threadfence();
    return value;
}

__global__ void initAtomic(unsigned& count, const unsigned initValue) {
    count = initValue;
}

__global__ void addVerify(unsigned& count, const unsigned biasAtomicValue) {
    atomicAdd(&count, 1);
    // NOTE: When uncomment the following while loop the addVerify is stuck,
    //       it cannot read last proper value in variable count
//    while (atomicLoad(&count) != (1024 * 1024 + biasAtomicValue)) {
//        printf("count = %u\n", atomicLoad(&count));
//    }
}

int main() {
    std::cout << "Hello, CUDA atomics!" << std::endl;
    const auto atomicSize = sizeof(unsigned);

    unsigned* datomic = nullptr;
    cudaMalloc(&datomic, atomicSize);

    cudaStream_t stream;
    cudaStreamCreate(&stream);

    constexpr unsigned biasAtomicValue = 11;
    initAtomic<<<1, 1, 0, stream>>>(*datomic, biasAtomicValue);
    addVerify<<<1024, 1024, 0, stream>>>(*datomic, biasAtomicValue);
    cudaStreamSynchronize(stream);

    unsigned countHost = 0;
    cudaMemcpyAsync(&countHost, datomic, atomicSize, cudaMemcpyDeviceToHost, stream);
    assert(countHost == 1024 * 1024 + biasAtomicValue);

    cudaStreamDestroy(stream);

    return 0;
}

セクションのコメントをatomicLoadで解除すると、アプリケーションが動かなくなります...

多分私は何かを逃した？アトミックに変更された変数をロードする適切な方法はありますか?

PS: 実装が存在することは知っていcuda::atomicますが、この API は私のハードウェアではサポートされていません

4

1 に答える 1

4

ワープはロックステップ方式で (少なくとも古いアーキテクチャでは) 動作するため、あるスレッドの条件付き待機と別のスレッドのプロデューサーを同じワープに置くと、ワープが開始/最初に実行されます。おそらく、非同期ワープスレッドスケジューリングを備えた最新のアーキテクチャのみがこれを行うことができます. たとえば、これを実行する前に、cuda アーキテクチャのマイナー バージョンとメジャー バージョンを照会する必要があります。Volta以降はOKです。

また、100 万のスレッドを起動し、それらすべてを一度に待機しています。GPU には、実行中の 100 万のスレッドを保持するための実行ポート/パイプラインの可用性があまりない場合があります。おそらく、64k CUDA パイプラインの GPU でのみ機能するでしょう (パイプラインごとに 16 スレッドが実行中であると仮定)。何百万ものスレッドを待機する代わりに、条件が発生したときにメイン カーネルからサブカーネルを生成するだけです。動的並列処理が重要な機能です。また、誰かが古代の nvidia カードを使用している場合に備えて、動的並列処理を使用するための最小のマイナー/メジャー cuda バージョンを確認する必要があります。

Atomic-add コマンドは、ターゲット アドレスの古い値を返します。条件の後で 3 番目のカーネルを 1 回だけ呼び出すつもりである場合は、動的並列処理を開始する前に、その戻り値を「if」で簡単に確認できます。

100 万回印刷していますが、パフォーマンスが良くなく、CPU/RAM が遅い場合、コンソール出力にテキストが表示されるまでに時間がかかる場合があります。

最後に、アトミック操作のパフォーマンスを最適化するには、最初に共有メモリで実行してから、ブロックごとに 1 回だけグローバル アトミックに移行します。これは、条件値よりも多くのスレッドがある場合 (常に 1 つのインクリメント値を想定)、条件のポイントを見逃すため、すべてのアルゴリズムに適用できるとは限りません。

于 2022-02-05T11:31:07.550 に答える