atomicInc() ne fonctionne pas

Question

J'ai essayé ci-dessous le programme en utilisant atomicInc().

__global__ void ker(int *count) 
{ 
    int n=1; 
    int x = atomicInc ((unsigned int *)&count[0],n); 
    CUPRINTF("In kernel count is %d\n",count[0]); 
} 

int main() 
{ 
    int hitCount[1]; 
    int *hitCount_d; 

    hitCount[0]=1; 
    cudaMalloc((void **)&hitCount_d,1*sizeof(int)); 

    cudaMemcpy(&hitCount_d[0],&hitCount[0],1*sizeof(int),cudaMemcpyHostToDevice); 

    ker<<<1,4>>>(hitCount_d); 

    cudaMemcpy(&hitCount[0],&hitCount_d[0],1*sizeof(int),cudaMemcpyDeviceToHost); 

    printf("count is %d\n",hitCount[0]); 
    return 0; 
} 

sortie est:

In kernel count is 1 
In kernel count is 1 
In kernel count is 1 
In kernel count is 1 

count is 1 

Je ne comprends pas pourquoi il n'incrémente. Quelqu'un peut-il aider à

Answer 1

Se référant à la documentation, atomicInc fait ceci:

pour ce qui suit:

atomicInc ((unsigned int *)&count[0],n); 

Compute:

((count[0] >= n) ? 0 : (count[0]+1)) 

et stocker le résultat de retour dans count[0]

(Si vous n'êtes pas su re ce que l'opérateur ? fait, regardez here)

Depuis que vous avez passé n = 1 et count[0] commence à 1, atomicInc jamais incrémente en fait la count[0] au-delà de la variable 1.

Si vous voulez le voir incrémenter au-delà de 1, passer une valeur plus grande pour n.

La variable n agit en fait comme une "valeur de survol" pour le processus d'incrémentation. Lorsque la variable à incrémenter atteint en réalité la valeur n, le atomicInc le réinitialise à zéro.

Bien que vous n'ayez pas posé la question, vous pourriez demander: «Pourquoi est-ce que je ne vois jamais une valeur de zéro si j'atteins la valeur de roulement?

Pour répondre à cette question, vous devez vous rappeler que tous les 4 de vos threads sont en cours d'exécution. Tous les 4 exécutent l'instruction atomicInc avant d'exécuter l'instruction d'impression suivante.

Nous avons donc une variable de count[0] qui commence à 1.

1. Le premier thread pour exécuter les réinitialisations atomiques à zéro.
2. Le thread suivant l'incrémente à 1.
3. Le troisième thread le réinitialise à zéro.
4. La quatrième et dernière par incréments de fil il à 1.

Ensuite, tous les 4 fils imprimer la valeur. Comme autre expérience, essayez de lancer 5 threads au lieu de 4, voir si vous pouvez prédire quelle sera la valeur imprimée.

ker<<<1,5>>>(hitCount_d); 

Comme @talonmies indiqué dans les commentaires, si vous remplacez votre atomicInc pour un atomicAdd:

int x = atomicAdd ((unsigned int *)&count[0],n); 

Vous obtiendrez des résultats que vous attendiez probablement.