atomicAddを使用してCUDAで画像を平均化すると、結果が一貫しない

Question

大きな問題では単純なテストケースとしてイメージの強度を平均化しています。しかし、私が得る結果は、それを実行するたびに少しずつ異なります。逆に、同じアルゴリズムをCPU上で連続して実行すると、結果は静的になります。のは、GPU上のコードを見てみましょう、

//util.cu 
__global__ void avgImageDevice(float3 *avg, float3 *d_colorImageRGB, unsigned int width, unsigned int height) 
{ 
    const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; 
    const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; 
    if (x >= width || y >= height) return; 

    atomicAdd(&avg->x, d_colorImageRGB[y*width + x].x); 
    atomicAdd(&avg->y, d_colorImageRGB[y*width + x].y); 
    atomicAdd(&avg->z, d_colorImageRGB[y*width + x].z); 
} 

extern "C" void avgImage(float3 *avg, float3 *d_colorImageRGB, unsigned int width, unsigned int height) 
{ 
    const int T_PER_BLOCK = 16; 
    const dim3 blockSize((width + T_PER_BLOCK - 1)/T_PER_BLOCK, (height + T_PER_BLOCK - 1)/T_PER_BLOCK); 
    const dim3 gridSize(T_PER_BLOCK, T_PER_BLOCK); 

    avgImageDevice << <blockSize, gridSize >> >(avg, d_colorImageRGB, width, height); 
} 

そして、CPUの実装は

//main.cpp 

#include <vector_types.h> 
#include <opencv2\core\core.hpp> 
#include <cuda_runtime.h> 
#include <string> 
extern "C" void avgImage(float3 *avg, float3 *d_colorImageRGB, unsigned int width, unsigned int height); 

int main() 
{ 
    for(int k = 0 ; k < 100 ;++k) 
    { 
     //Initialization 
     Mat Image; 
     float3 avgCPU = make_float3(0, 0, 0); 
     float3 avgGPU = make_float3(0, 0, 0); 
     std::string filenameImage("/foo.jpg"); 
     Image = imread(filenameImage, -1); 
     Image.convertTo(Image, CV_32FC3, 1.0f/255); 

     //Copy to GPU global memory 
     cutilSafeCall(cudaMemcpy(d_albedoMapFilteredFloat3, Image.data, sizeof(float) * 3 * Image.size().width * Image.size().height, cudaMemcpyHostToDevice)); 

     //Average on CPU 
     for (int x = 0; x < Image.size().width; ++x) 
      for (int y = 0; y < Image.size().height; ++y) 
      { 
       Vec3f intensity = Image.at<Vec3f>(y, x); 
       avgCPU += make_float3(intensity.val[0], intensity.val[1], intensity.val[2]); 
      } 
     avgCPU /= Image.size().width * Image.size().height; 


     //Average on GPU 
     float3 *d_avg; 
     cutilSafeCall(cudaMalloc(&d_avg, sizeof(float3))); 
     cutilSafeCall(cudaMemset(d_avg, 0, sizeof(float3))); 
     avgImage(d_avg, d_albedoMapFilteredFloat3, Image.size().width, Image.size().height); 
     cutilSafeCall(cudaMemcpy(&avgGPU, d_avg, sizeof(float3), cudaMemcpyDeviceToHost)); 
     avgGPU /= Image.size().width * Image.size().height; 

     //Following values are consant across the iterations 
     printf("AVG CPU r: %.10f, g: %.10f, b: %.10f\n", avgCPU.x, avgCPU.y, avgCPU.z); 

     //Following values are different at every iteration 
     printf("AVG GPU r: %.10f, g: %.10f, b: %.10f\n", avgGPU.x, avgGPU.y, avgGPU.z); 
    } 
} 

だから、次の行の各ペアが一致し、かつ静的である必要があり、次の通りです。しかし、それらは一致しないし、GPUの結果は静的ではない。

AVG CPU r: 0.6326226592, g: 0.6762236953, b: 0.6836426258 
AVG GPU r: 0.6325752139, g: 0.6762712002, b: 0.6835504174 
AVG CPU r: 0.6326226592, g: 0.6762236953, b: 0.6836426258 
AVG GPU r: 0.6325753927, g: 0.6762660146, b: 0.6835544705 
AVG CPU r: 0.6326226592, g: 0.6762236953, b: 0.6836426258 
AVG GPU r: 0.6325772405, g: 0.6762678027, b: 0.6835457087 
AVG CPU r: 0.6326226592, g: 0.6762236953, b: 0.6836426258 
AVG GPU r: 0.6325744987, g: 0.6762621403, b: 0.6835452914 
AVG CPU r: 0.6326226592, g: 0.6762236953, b: 0.6836426258 
AVG GPU r: 0.6325761080, g: 0.6762756109, b: 0.6835403442 
AVG CPU r: 0.6326226592, g: 0.6762236953, b: 0.6836426258 
AVG GPU r: 0.6325756311, g: 0.6762655973, b: 0.6835408211 

私はGTX 960、CUDA 6.5、およびWindows 7を使用しています。これはデータ競合の問題ですか？私の知る限り、atomicAddはグローバルメモリに問題がないと報告されています。

Answer 1

これはデータ競争ではありません。

浮動小数点の加算は可換である：

a + b == b + a 

しかし、それはない連想です。

(a + b) + c != a + (b + c) 

（特にそれらが関連付けられている方法）は、個々の追加の異なる順序が異なる結果を与える：A、B、Cのように存在します。

Answer 2

結果は、スレッドのスケジューリング順序によって異なる場合があります。実際には、イメージのサイズ、コンポーネントの値によって、結果の平均値は実行ごとに多少異なる場合がありますが、すべての値が正しいとは限りません。ランが違う場合は、コードの他の部分に問題がある可能性が最も高いです。 numが同じ場合、すべての結果はIEEE-754ノルムまで正しくなります。