OpenCLで浮動小数のatomic addをしたい

atomic addとは、あるメモリ領域に複数スレッドから数値を加えたい時に、アクセスのコンフリクトを避けるために排他制御を行い正しく和を計算するための処理のこと。

CUDAでdoubleのatomicAddを行う方法はCUDA C Programming Guideに書いてある。CUDA Toolkit Documentationを参照。
コードをここに引用する。

__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)));
  } while (assumed != old);
  return __longlong_as_double(old);
}

atomicCASとはcompare and swapのことらしい。

ではOpenCLではどうしたら良いか。そのまんま書き換えて大丈夫だった。
atom_cmpxchgというものを使えばよい。おそらくcompare exchangeの意かと。

floatの場合

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

float atom_add_float(__global float* const address, const float value)
{
  uint oldval, newval, readback;
  
  *(float*)&oldval = *address;
  *(float*)&newval = (*(float*)&oldval + value);
  while ((readback = atom_cmpxchg((__global uint*)address, oldval, newval)) != oldval) {
    oldval = readback;
    *(float*)&newval = (*(float*)&oldval + value);
  }
  return *(float*)&oldval;
}


doubleの場合

#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable

double atom_add_double(__global double* const address, const double value)
{
  long oldval, newval, readback;
  
  *(double*)&oldval = *address;
  *(double*)&newval = (*(double*)&oldval + value);
  while ((readback = atom_cmpxchg((__global long*)address, oldval, newval)) != oldval) {
    oldval = readback;
    *(double*)&newval = (*(double*)&oldval + value);
  }
  return *(double*)&oldval;
}

OpenCLは関数名のOverLoadが出来ないのかな?いちいち名前を変えないといけない…CUDAの場合は関数名の前に__OVERLOADABLE__修飾子を付けたらビルド出来たけれど、Appleだとだめだった。

試しにTesla C2050でdoubleの1.0を10万個ほどatom_addしてみたが、5秒程度かかった。


単純に総和計算をしたい場合は、実装はやや面倒だが

の方法を使えば計算は断然速くなる。
atom_addはどうしても必要な時以外は避けたほうが良い。