Étalonner une image sur CUDA en utilisant atomicAdd donne des résultats incohérents

Question

Je fais la moyenne de l'intensité d'une image comme un cas de test simple pour un problème plus important. Mais le résultat que je reçois est légèrement différent chaque fois que je le lance. Inversement, si j'exécute séquentiellement le même algorithme sur la CPU, le résultat est statique. Regardons le code sur 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); 
} 

Et la mise en œuvre du processeur est la suivante,

//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); 
    } 
} 

Ainsi, chaque paire de lignes suivantes doivent correspondre et être statique. Mais ils ne correspondent pas et les résultats du GPU ne sont pas statiques.

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 

J'ai une GTX 960, CUDA 6.5 et Windows 7. Est-ce une question course de données? Au meilleur de ma connaissance, atomicAdd n'a pas signalé de problème avec la mémoire globale.

Answer 1

Ce n'est pas une course de données.

plus de virgule flottante est commutative:

a + b == b + a 

Mais il est pas associative; il y a, b, c tel que:

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

Les différents ordres des additions individuelles (en particulier la façon dont ils sont associés) donnera des résultats différents.

Answer 2

Le résultat peut dépendre de l'ordre de programmation des threads. En effet, en fonction de la taille de l'image, des valeurs dans les composants, la moyenne obtenue peut légèrement différer d'une exécution à l'autre, bien que toutes les valeurs soient correctes. Si num diffère d'une exécution à une autre, il y a probablement un problème dans d'autres parties du code. Si num est le même, tous vos résultats sont corrects jusqu'à la norme IEEE-754.