CUDA - atomicAdd（float）不會添加非常小的值

當我使用float atomicAdd(float *address, float val)添加小於約的浮點值時。 1e-39到0，加不工作，並在address值保持爲0CUDA - 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：只有一個線程在運行。編譯器在第二種情況下不會將其清除爲0，因此在第一種情況下沒有理由將其清除。此外，這是一個最小的代碼示例，它是由複雜的動態數值集成的結果分配了更復雜的代碼片段所致，因此編譯器無法刷新它。 – user2412789

@ sgar91：爲了防止可能的編譯器優化，我將賦值添加爲addit = sinf（value）而不是addit = value。 – user2412789

atomicAdd是一種特殊的指令並不一定遵循相同的沖洗和舍入的行爲，如果你指定例如-ftz=true或-ftz=false其他浮點運算（如普通FP加）

正如記載，你可能會得到PTX ISA manual：

浮點運算.add是一個單精度32位運算。 atom.add.f32舍入到最接近的偶數，並刷新低於正常的輸入和結果以保留零。

因此，即使普通的浮點加法應該如果指定-ftz=false不沖水非正規數爲零（這是默認的，我相信，對於nvcc），浮點原子加操作全局內存將刷新到零（總是）。

來源

2013-11-25 18:07:22

CUDA - atomicAdd（float）不會添加非常小的值

回答

相關問題