當我使用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。
'1e-40'是單精度浮點格式的非正規數,並且超出'float'數據類型的精度範圍。編譯器很可能會將其刷新爲零。另外,如果內核在多線程上啓動,那麼在同一內存位置上執行'+ ='操作將導致未定義的行爲。 – sgarizvi
@ sgar91:只有一個線程在運行。編譯器在第二種情況下不會將其清除爲0,因此在第一種情況下沒有理由將其清除。此外,這是一個最小的代碼示例,它是由複雜的動態數值集成的結果分配了更復雜的代碼片段所致,因此編譯器無法刷新它。 – user2412789
@ sgar91:爲了防止可能的編譯器優化,我將賦值添加爲addit = sinf(value)而不是addit = value。 – user2412789