原子操作与double，OpenCL [英] Atomic operations with double, OpenCL

问题描述

我想知道是否有一种实现原型操作（特别是atomic_add）的双重类型。

I would like to know if there's a way to implement atomic operations (particularly atomic_add) with double type.

对于浮点数，这段代码可以工作，但是 atomic_xchg 不支持double：

For floats this code works, but atomic_xchg doesn't support double:

while ((value = atomic_xchg(addr, atomic_xchg(addr, 0.0f)+value))!=0.0f);

推荐答案

过去我一直在寻找，我发现这个： https://github.com /ddemidov/vexcl-experiments/blob/master/sort-by-key-atomic.cpp 。
最后我想出了不同的方法来解决我的问题，所以我没有使用它。这是代码：

I was looking for for the same in the past and I found this: https://github.com/ddemidov/vexcl-experiments/blob/master/sort-by-key-atomic.cpp. At the end I figured out different approach to my problem so I did not use it. Here is the code:

    "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n"
    "#pragma OPENCL EXTENSION cl_khr_int64_base_atomics: enable\n"
    "void AtomicAdd(__global double *val, double delta) {\n"
    "  union {\n"
    "    double f;\n"
    "    ulong  i;\n"
    "  } old;\n"
    "  union {\n"
    "    double f;\n"
    "    ulong  i;\n"
    "  } new;\n"
    "  do {\n"
    "    old.f = *val;\n"
    "    new.f = old.f + delta;\n"
    "  } while (atom_cmpxchg ( (volatile __global ulong *)val, old.i, new.i) != old.i);\n"
    "}\n"
    "kernel void atomic_reduce(\n"
    "  ulong n,\n"
    "  global const int    * key,\n"
    "  global const double * val,\n"
    "  global double * sum\n"
    ")\n"
    "{\n"
    "  for(size_t idx = get_global_id(0); idx < n; idx += get_global_size(0))\n"
    "    AtomicAdd(sum + key[idx], val[idx]);\n"
    "}\n",
    "atomic_reduce"

这篇关于原子操作与double，OpenCL的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！