La fonction du noyau CUDA semble montrer des conditions de course malgré des conditions de course

Question

Ma fonction noyau CUDA ne retourne pas le résultat attendu (une somme de tous les éléments du vecteur b) mais retourne une seule valeur du vecteur b. J'ai essayé memcheck et racecheck, mais rien venu:

[breecej@compute-0-32 newsum]$ cuda-memcheck mystock 
========= CUDA-MEMCHECK 
========= ERROR SUMMARY: 0 errors 
[breecej@compute-0-32 newsum]$ cuda-memcheck --tool racecheck mystock 
========= CUDA-MEMCHECK 
========= RACECHECK SUMMARY: 0 hazards displayed (0 errors, 0 warnings) 
[breecej@compute-0-32 newsum]$ 

Voici la fonction du noyau:

__global__ void AddDoubles(double *a, double *b, int count) 
{ 
    int id = blockIdx.x * blockDim.x + threadIdx.x; 
    if(id < count) 
    { 
     a[0] += b[id]; 
     __syncthreads(); 
    } 
} 

Et voici les variables d'entrer dans et la façon dont je l'appeler dans le principal:

int count = vect.size(); 
//allocate memory for the sum and stock values for the host side 
double *h_a = new double[1]; 
double *h_b = new double[count]; 
//array a set to 0, array b set to the values from the line 
h_a[0] = 0; 
for(int i = 0; i < count; i++) 
{ 
    h_b[i] = vect.at(i); 
} 
//allocate memory for the sum and stock values arrays for the device side 
double *d_a, *d_b; 
if(cudaMalloc(&d_a, sizeof(double)) != cudaSuccess) 
{ 
    cout << "Nope! a did not allocate correctly"; 
    return 0; 
} 
if(cudaMalloc(&d_b, sizeof(double) * count) != cudaSuccess) 
{ 
    cout << "Nope! b did not allocate correctly"; 
    cudaFree(d_a); 
    return 0; 
} 
//copy the host pointer to the device pointer 
if(cudaMemcpy(d_a, h_a, sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess) 
{ 
    cout << "Could not copy!" << endl; 
    cudaFree(d_a); 
    cudaFree(d_b); 
    return 0; 
} 
if(cudaMemcpy(d_b, h_b, sizeof(double) * count, cudaMemcpyHostToDevice) != cudaSuccess) 
{ 
    cout << "Could not copy!" << endl; 
    cudaFree(d_a); 
    cudaFree(d_b); 
    return 0; 
} 
//use AddDoubles to sum up all of the values in b and put them into a 
AddDoubles<<<count/256 + 1, 256>>>(d_a, d_b, count); 

où "vect" est un vecteur de doubles. Pourquoi semble-t-il qu'il existe une course pour changer la valeur dans la fonction du noyau, mais rien ne se présente dans racecheck?

Answer 1

Votre code a une course mémoire globale, mais cuda-memcheck ne le détectera pas. Citant le documentation:

Le racecheck outil est un détecteur de danger pour l'accès aux données de la mémoire partagée de temps d'exécution . L'utilisation principale de cet outil est d'aider à identifier la mémoire conditions de course d'accès dans les applications CUDA qui utilisent la mémoire partagée.

cuda-memcheck ne peut détecter que les courses de mémoire partagée. Votre code n'utilise pas la mémoire partagée.

La course elle-même peut être fait évident si votre noyau est écrit comme ceci:

__global__ void AddDoubles(double *a, double *b, int count) 
{ 
    int id = blockIdx.x * blockDim.x + threadIdx.x; 
    if(id < count) 
    { 
     double x = a[0]; // 1. load a[0] to thread local register 
     double y = b[id]; // 2. load b[id] to thread local register 
     double z = x + y; // 3. perform addition in thread local register 
     a[0] = z;   // 4. store thread local register sum to a[0] 
    } 
} 

Cela ne peut jamais être correcte si l'exécution est sérialisé. Si un thread stocke à a[0] alors qu'un autre thread se trouve entre les étapes 1 et 4, le contenu de a[0] sera invalidé par la deuxième écriture. Dans un modèle d'exécution pipeliné massivement parallèle comme CUDA, cela va de soi.

Notez également que votre utilisation de __syncthreads() n'a aucun effet sur ce comportement, et le noyau fonctionnera de manière identique avec ou sans son inclusion dans le code. Pour comprendre comment effectuer ce type de réduction en parallèle dans CUDA, reportez-vous au CUDA reduction example, qui comprend un excellent livre blanc sur les options de fonctionnement et d'optimisation des performances.