Coder un noyau CUDA dont plusieurs threads écrivent dans le même index?

Question

Je suis en train d'écrire du code pour activer les réseaux de neurones sur CUDA, et je suis confronté à un problème. Je ne reçois pas la somme correcte des poids entrant dans un neurone donné.

Donc, voici le code du noyau, et je vais essayer de l'expliquer un peu plus clairement avec les variables.

__global__ void kernelSumWeights(float* sumArray, float* weightArray, int2* sourceTargetArray, int cLength) 
{ 
int nx = threadIdx.x + TILE_WIDTH*threadIdx.y; 
int index_in = (blockIdx.x + gridDim.x*blockIdx.y)*TILE_WIDTH*TILE_WIDTH + nx; 
if(index_in < cLength) 
{ 

    sumArray[sourceTargetArray[index_in].y] += fabs(weightArray[index_in]); 
    //__threadfence(); 
    __threadfence_block(); 

} 

} 

Tout d'abord, le nombre de connexions dans le réseau est cLength. Pour chaque connexion, il y a un neurone source et un neurone cible, ainsi qu'un poids pour cette connexion. SourceTargetArray contient cette information. Ainsi, l'index i de sourceTargetArray est l'index de neurones source de la connexion i, et l'index neuronal cible de la connexion i. Le weightArray contient les informations de poids (l'index i de weightArray correspond à la connexion i).

Comme vous pouvez le voir, SumArray est l'endroit où je stocke les sommes. Ainsi, le noyau incrémente le sumArray (à l'index de neurones cible de la connexion i) par la valeur absolue du poids de la connexion i. Intuitivement, pour toutes les connexions entrantes au neurone, additionnez tous les poids. C'est vraiment tout ce que j'essaie de faire avec ce noyau. Finalement, je vais normaliser les poids en utilisant cette somme.

Le problème est que c'est faux. Je l'ai fait en série, et la réponse est différente. La réponse diffère, généralement d'environ 12-15x (donc la bonne réponse sera 700.0 et ce que je reçois est quelque chose dans la gamme des années 50).

Vous pouvez voir que j'ai ajouté __threadfence() (et __threadfence_block() pour essayer de m'assurer que les écritures n'étaient pas effectuées en même temps par chaque thread). Je ne suis pas sûr si c'est le problème avec mon code. Je me suis assuré que le tableau de poids est identique à la version en série que j'ai testé, et que les informations source/cible sont identiques. Qu'est-ce que je fais mal?

EDIT: Pour référence, __threadfence() usaged est décrite dans le Guide de programmation CUDA v3.1 Annexe B.5 Fonctions de clôture de la mémoire

Answer 1

Vous devez faire une réduction.

Somme des éléments affectés à chaque fil et placer le résultat dans un tableau, cache [threadsPerBlock] puis __Syncthreads

maintenant réduire les totaux sous résultant par addition de sous-totaux voisins successifs:

int cacheIndex = threadIdx.x; 
int i = blockDim.x/2; 
while (i != 0) 
{ 
    if (cacheIndex < i) 
     cache[cacheIndex] += cache[cacheIndex] + 1; 
     __syncthreads; 
     i /= 2; 
    } 
} 

Ci-dessous pont explique cette question en détail:

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

Exemple de code pour c'est son e:

http://www.nvidia.com/object/cuda_sample_data-parallel.html

Il est également très bien expliqué dans "CUDA par l'exemple" (ce qui est l'endroit où le fragment de code provient).

Il y a une grande mise en garde avec cette approche. Les ajouts ne se produiront pas dans le même ordre qu'ils le feraient avec le code série. L'ajout de flottants n'est pas commutatif, de sorte que des erreurs d'arrondi peuvent conduire à des résultats légèrement différents.

Answer 2

+= n'est pas atomistique => pas thread-safe. Utilisez atomicAdd.

Vous devriez également éviter d'écrire dans la même cellule mémoire. Le problème est que ces appels seront sérialisés, les threads feront la queue et attendront l'un l'autre. Si vous ne pouvez pas éviter cette opération, essayez de décomposer votre algorithme en deux phases: le calcul individuel et la fusion. La fusion parallèle peut être implémentée très efficacement.