c++ - 在 openCL 中的非整数上实现 atom_add 函数，具有较旧计算能力的设备

Question

我想在具有非整数(float 和 double)的设备内存上使用原子函数，例如我在 CUDA C 编程指南 中看到了实现 atomicAdd< 的下一个代码 double float 函数:

从 CUDA C Programming Guide 中提取的代码:

---

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));

    // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif

是否可以在 openCL 中做类似的事情？我有一台计算能力为 2.1 的设备

UPD

---

我设法编写了一个似乎有效的与原始代码等效的代码:

double atom_add_double(__global double* address, double val) {
    __global long* address_as_ull =
        (__global long*)address;
    long old = *address_as_ull;
    long assumed;

    do {
        assumed = old;
        old = atom_cmpxchg(address_as_ull, assumed,
            as_long(val + as_double(assumed)));
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
    } while (assumed != old);

    return as_double(old);
}

帖子的回复中有更多详细信息，感谢@pmdj。

Answer 1

对于 64 位数据 (double)，您需要测试 cl_khr_int64_base_atomics扩大。如果您的实现支持，您可以将 atom_cmpxchg() 函数与 long/ulong(64 位整数)值一起使用。

对于 32 位 float，atomic_cmpxchg function是核心 OpenCL 1.2 及更高版本的一部分。如果您的实现仅支持 OpenCL 1.0，则需要测试 the cl_khr_global_int32_base_atomics extension并在支持的情况下使用其 atom_cmpxchg() 函数。

在 OpenCL 中使用 as_typen 运算符最容易将浮点值的二进制表示视为整数，反之亦然。 (或者，您可以为此使用 union 类型，尽管在这种情况下这没有帮助；有关详细信息，请参阅 OpenCL 规范的第 6.2.4 节。)在您的代码中，等同于 __double_as_longlong 将是 as_long()，而不是 __longlong_as_double 您将使用 as_double()。