Почему AtomAdd не был реализован для парных?
Почему hasnt atomicAdd() для двойников был реализован явно как часть CUDA 4.0 или выше?
В приложении F Страница 97 Руководства по программированию CUDA 4.1 приведены следующие версии
atomicAdd реализованы.
int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
unsigned long long int val);
float atomicAdd(float* address, float val)
На этой же странице приводится небольшая реализация atomicAdd для двойников следующим образом
которые я только начал использовать в своем проекте.
__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);
}
Почему бы не определить вышеуказанный код как часть CUDA?
Ответы
Ответ 1
Изменить: с CUDA 8 в CUDA реализована двухточечная точность atomicAdd() с аппаратной поддержкой графических процессоров SM_6X (Pascal).
В настоящее время никакие устройства CUDA не поддерживают atomicAdd для double в аппаратном обеспечении. Как вы отметили, он может быть реализован в терминах atomicCAS для 64-битных целых чисел, но есть нетривиальные эксплуатационные расходы для этого.
Следовательно, команда программного обеспечения CUDA решила задокументировать правильную реализацию как вариант для разработчиков, а не сделать ее частью стандартной библиотеки CUDA. Таким образом, разработчики не бессознательно выбирают стоимость исполнения, которую они не понимают.
Кроме того: я не думаю, что этот вопрос следует закрыть как "не конструктивный". Я думаю, что это совершенно правильный вопрос, +1.
