2

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

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

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

4

1 回答 1

4

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

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

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

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

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

于 2022-02-05T11:31:07.550 回答