时间戳

首页 > 解决方案 > 在 openCL 中的非整数上实现 atom_add 函数，具有较旧计算能力的设备

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

问题描述

我想在具有非整数（浮点数和双精度数）的设备内存上使用原子函数，例如，我在CUDA C Programming GuideatomicAdd中看到了双精度浮点数的实现函数的下一个代码：

从CUDA C 编程指南中提取的代码：

#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，在帖子的回复中提供了更多详细信息。

标签： c++opencl

解决方案

对于 64 位数据 ( double)，您需要测试cl_khr_int64_base_atomics扩展。如果您的实现支持这一点，您可以使用atom_cmpxchg()带有long/ ulong（64 位整数）值的函数。

对于 32 位float，该atomic_cmpxchg函数是核心 OpenCL 1.2 及更高版本的一部分。如果您的实现仅支持 OpenCL 1.0，您将需要测试扩展并使用其功能cl_khr_global_int32_base_atomicsatom_cmpxchg()（如果支持）。

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

推荐阅读