是否有适当的 CUDA atomicLoad 功能？

Question

我遇到过 CUDA atomic 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 取消注释该部分，则应用程序将卡住...

也许我错过了什么？有没有合适的方法来加载原子修改的变量？

P.S.：我知道有 cuda::atomic 实现，但是我的硬件不支持这个 API

Answer 1

由于 warp 以锁步方式工作（至少在旧拱门中），如果您将一个线程和一个生产者放在另一个线程上的条件等待，两者都在同一个 warp 中，那么这个 warp 可能会陷入等待状态首先开始/执行。也许只有具有异步扭曲线程调度的最新架构才能做到这一点。例如，您应该在运行之前查询 cuda 架构的次要版本。 Volta 及以上版本都可以。

此外，您正在启动 100 万个线程并同时等待所有线程。 GPU 可能没有那么多执行端口/管道可用性来运行 100 万个线程。也许它只能在 64k CUDA 管道的 GPU 中工作（假设每个管道有 16 个线程在运行）。无需等待数百万个线程，只需在条件发生时从主内核生成子内核。动态并行是关键特性。您还应该检查最小的次要 cuda 版本以使用动态并行，以防有人使用古老的 nvidia 卡。

Atomic-add 命令返回目标地址中的旧值。如果您打算仅在条件之后调用第三个内核一次，那么您可以在开始动态并行之前通过“if”简单地检查返回值。

您正在打印 100 万次，这对性能不利，如果您的 CPU/RAM 较慢，可能需要一些时间才能在控制台输出中显示文本。

最后，您可以优化原子操作的性能，方法是先在共享内存上运行它们，然后每个块只运行一次全局原子操作。如果线程数多于条件值（假设始终为 1 个增量值），这将错过条件点，因此它可能不适用于所有算法。