CUDA - atomicAdd (float) не добавляет очень малых значений

Когда я использую float atomicAdd(float *address, float val), чтобы добавить значение поплавка меньше, чем ок. 1e-39 к 0, добавление не работает, а значение при address остается 0.CUDA - atomicAdd (float) не добавляет очень малых значений

Вот простейший код:

__device__ float test[6] = {0}; 
__global__ void testKernel() { 
    float addit = sinf(1e-20); 
    atomicAdd(&test[0], addit); 
    test[1] += addit; 
    addit = sinf(1e-37); 
    atomicAdd(&test[2], addit); 
    test[3] += addit; 
    addit = sinf(1e-40); 
    atomicAdd(&test[4], addit); 
    test[5] += addit; 
}

Когда я запускаю код, указанный выше в качестве testKernel<<<1, 1>>>(); и остановить с помощью отладчика я вижу :

test 0x42697800 
    [0] 9.9999997e-21 
    [1] 9.9999997e-21 
    [2] 9.9999999e-38 
    [3] 9.9999999e-38 
    [4] 0    
    [5] 9.9999461e-41

Обратите внимание на разницу между тестом [4] и тестом [5]. Оба делали то же самое, но простое дополнение работало, и атомный ничего не делал. Что мне здесь не хватает?

Обновление: Системная информация: CUDA 5.5.20, карта NVidia Titan, драйвер 331.82, Windows 7x64, Nsight 3.2.1.13309.

источник

2013-11-25 user2412789

'1e-40' - это денормальное число для формата с плавающей запятой с одинарной точностью и выходит за пределы прецизионного диапазона типа данных« float ». Скорее всего, компилятор сбрасывает его до нуля. Кроме того, если ядро запускается в нескольких потоках, операция '+ =' в том же месте памяти приведет к неопределенному поведению. – sgarizvi

@ sgar91: Есть только 1 поток. Компилятор не очищает его до нуля во втором случае, поэтому нет причин для его сброса в первом случае. Кроме того, это минимальный пример кода, который был мотивирован гораздо более сложным фрагментом кода, в котором addit был назначен в результате сложного динамического численного интегрирования, поэтому компилятор не смог его очистить. – user2412789

@ sgar91: Чтобы предотвратить возможную оптимизацию компилятора, я изменил назначение на addit на addit = sinf (значение) вместо addit = value. – user2412789

atomicAdd специальная инструкция, которая не обязательно выполнять тот же флеш и округление поведения, которые вы могли бы получить, если указать, например, -ftz=true или -ftz=false на других операциях с плавающей точкой (например, обычный Ф.П. надстройка)

Как указано в PTX ISA manual:

Операция с плавающей запятой .add - это 32-разрядная операция с одной точностью. atom.add.f32 до ближайшего четного и сбрасывает субнормальные входы и результаты для сохранения нуля нуля.

Так что даже если обычный с плавающей запятой надстройкой не должны подрозетниками денормализованных чисел к нулю, если вы укажете -ftz=false (который по умолчанию, я считаю, для nvcc), с плавающей точкой атомного операцию добавления к глобальной памяти смоет к ноль (всегда).

источник

2013-11-25 18:07:22

CUDA - atomicAdd (float) не добавляет очень малых значений

ответ

Смежные вопросы