реализация функции atom_add на нецелых числах в openCL, устройствах с более старыми вычислительными возможностями

Question

Я хочу использовать элементарные функции в памяти устройства с нецелыми числами (с плавающей и двойной), например, я видел в CUDA C Руководство по программированию следующий код для реализации функции atomicAdd для двойной-точные числа с плавающей точкой:

код, извлеченный из Руководство по программированию в 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.

Answer 1

Для 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_atomics и использовать его функцию atom_cmpxchg(), если поддерживается.

Обработка двоичного представления с плавающей точкойЗначение в виде целого числа и наоборот наиболее легко делается в OpenCL с использованием оператора as_typen.(В качестве альтернативы вам разрешено использовать для этого типы union, хотя в данном случае это бесполезно; подробности см. В разделе 6.2.4 спецификации OpenCL.) В вашем коде эквивалент __double_as_longlong будет равен as_long()и вместо __longlong_as_double вы бы использовали as_double().