是否有适当的 CUDA atomicLoad 函数？

Question

我遇到过 CUDA atomic API 没有 atomicLoad 函数的问题。在 Whosebug 上搜索后，我发现了以下 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 以锁步方式工作（至少在旧 arch 中），如果您在同一个 warp 中对一个线程和生产者进行条件等待，那么 warp 可能会卡在等待中如果它 starts/is 先执行。也许只有具有异步 warp 线程调度的最新架构才能做到这一点。例如，您应该在运行之前查询 minor-major 版本的 cuda 体系结构。 Volta及以后的都可以。

您还启动了 100 万个线程并同时等待所有线程。 GPU 可能没有那么多执行 ports/pipeline 可用性来拥有 100 万个线程 in-flight。也许它只适用于 64k CUDA 管道的 GPU（假设每个管道有 16 个线程在运行）。无需等待数百万个线程，只需在条件发生时从主内核中生成 sub-kernels。动态并行是关键特性。您还应该检查最低 minor-major cuda 版本以使用动态并行性，以防有人使用古老的 nvidia 卡。

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

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

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

是否有适当的 CUDA atomicLoad 函数？

Is there proper CUDA atomicLoad function?

cuda

gpu-atomics