2010-09-22 5 views
5

Je calcule la distance euclidienne entre des points n-dimensionnels en utilisant OpenCL. Je reçois deux listes de points n-dimensionnels et je devrais retourner un tableau qui contient seulement les distances de chaque point de la première table à chaque point de la deuxième table.Récapitulatif des tableaux cumulés en utilisant OpenCL

Mon approche est de faire la boucle doble régulière (pour chaque point au tableau 1 {pour chaque point Tableau2 {...}} puis faire le calcul pour chaque paire de points dans paralell.

euclidienne la distance est ensuite divisée en 3 parties: 1. prendre la différence entre chaque dimension dans les points 2. mettre cette différence par carré (toujours pour chaque dimension) 3. additionner toutes les valeurs obtenues en 2. 4. Prendre la racine carrée de la valeur obtenue en 3. (cette étape a été omise dans cet exemple.)

Tout fonctionne comme un charme jusqu'à ce que je essayer d'accumuler la somme de toutes les différences (à savoir, exécuter l'étape 3 de la procédure décrite ci-dessus, ligne 49 du code ci-dessous). Comme données de test, j'utilise DescriptorLists avec 2 points chacune: DescriptorList1: 001,002,003, ..., 127,128; (p1) 129, 130, 131, ..., 255, 256; (p2)

DescriptorList2: 000,001,002, ..., 126,127; (p1) 128,129,130, ..., 254,255; (p2)

Donc le vecteur résultant devrait avoir les valeurs: 128, 2064512, 2130048, 128 En ce moment je reçois des nombres aléatoires qui varient avec chaque exécution.

J'apprécie toute aide ou conduit sur ce que je fais mal. Espérons que tout est clair sur le scénario que je travaille dans

#define BLOCK_SIZE 128 

typedef struct 
{ 
    //How large each point is 
    int length; 
    //How many points in every list 
    int num_elements; 
    //Pointer to the elements of the descriptor (stored as a raw array) 
    __global float *elements; 
} DescriptorList; 

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float As[BLOCK_SIZE]) 
{ 

    int gpidA = get_global_id(0); 

    int featA = get_local_id(0); 

    //temporary array to store the difference between each dimension of 2 points 
    float dif_acum[BLOCK_SIZE]; 

    //counter to track the iterations of the inner loop 
    int loop = 0; 

    //loop over all descriptors in A 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 

     //take the i-th descriptor. Returns a DescriptorList with just the i-th 
     //descriptor in DescriptorList A 
     DescriptorList tmpA = GetDescriptor(A, i); 

     //copy the current descriptor to local memory. 
     //returns one element of the only descriptor in DescriptorList tmpA 
     //and index featA 
     As[featA] = GetElement(tmpA, 0, featA); 
     //wait for all the threads to finish copying before continuing 
     barrier(CLK_LOCAL_MEM_FENCE); 

     //loop over all the descriptors in B 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      //take the difference of both current points 
      dif_acum[featA] = As[featA]-B.elements[k*BLOCK_SIZE + featA]; 
      //wait again 
      barrier(CLK_LOCAL_MEM_FENCE); 
      //square value of the difference in dif_acum and store in C 
      //which is where the results should be stored at the end. 
      C[loop] = 0; 
      C[loop] += dif_acum[featA]*dif_acum[featA]; 
      loop += 1; 
      barrier(CLK_LOCAL_MEM_FENCE); 
     } 
    } 
} 

Répondre

7

Votre problème réside dans ces lignes de code.

C[loop] = 0; 
C[loop] += dif_acum[featA]*dif_acum[featA]; 

Tous les sujets de votre groupe de travail (bien, en fait tous vos fils, mais Nous allons essayer de modifier cette position de tableau simultanément sans aucune synchronisation. Plusieurs facteurs font de ce vraiment problématique:

  1. Le groupe de travail est ne pas fonctionner complètement en parallèle, ce qui signifie que pour certains fils C [boucle] = 0 peut être appelé après d'autres discussions ont déjà exécuté la ligne suivante
  2. Ceux qui s'exécutent en parallèle lisent tous la même valeur à partir de C [boucle], la modifient avec leur incrément et essayent de réécrire à la même adresse. Je ne suis pas tout à fait sûr du résultat de cette réécriture (je pense que l'un des threads réussit à écrire, alors que les autres échouent, mais je ne suis pas complètement sûr), mais c'est faux de toute façon.

permet maintenant résoudre ce problème: Alors que nous pourrions être en mesure d'obtenir ce travail sur la mémoire globale à l'aide Atomics, il ne sera pas rapide, laisse accumuler dans la mémoire locale:

local float* accum; 
... 
accum[featA] = dif_acum[featA]*dif_acum[featA]; 
barrier(CLK_LOCAL_MEM_FENCE); 
for(unsigned int i = 1; i < BLOCKSIZE; i *= 2) 
{ 
    if ((featA % (2*i)) == 0) 
     accum[featA] += accum[featA + i]; 
    barrier(CLK_LOCAL_MEM_FENCE); 
} 
if(featA == 0) 
    C[loop] = accum[0]; 

de Bien sûr, vous pouvez réutiliser d'autres tampons locaux, mais je pense que le point est clair (btw: Êtes-vous sûr que dif_acum sera créé dans la mémoire locale, parce que je pense avoir lu quelque part que cela ne serait pas mis dans la mémoire locale rendrait le préchargement A dans la mémoire locale un peu inutile).

Quelques autres points sur ce code:

  1. Votre code est semble être adapté à l'utilisation que sur un groupe de travail (vous ne l'utilisez soit id groupId ni globale pour voir les éléments à travailler sur), pour performances optimales que vous pourriez vouloir utiliser plus que cela.
  2. est peut-être preferance personnelle, mais pour moi je, il semble préférable d'utiliser get_local_size(0) pour la workgroupsize que d'utiliser un define (puisque vous pouvez le changer dans le code hôte sans vous rendre compte devrait avoir changé votre code OpenCL à)
  3. Les barrières dans votre code sont toutes inutiles, car aucun thread n'accède à un élément de la mémoire locale qui est écrit par un autre thread. Vous n'avez donc pas besoin d'utiliser de mémoire locale pour cela.

Compte tenu de la dernière balle, vous pouvez simplement faire:

float As = GetElement(tmpA, 0, featA); 
... 
float dif_acum = As-B.elements[k*BLOCK_SIZE + featA]; 

Cela rendrait le code (ne pas considérer les deux premières balles):

__kernel void CompareDescriptors_deb(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE]) 
{ 
    int gpidA = get_global_id(0); 
    int featA = get_local_id(0); 
    int loop = 0; 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 
     DescriptorList tmpA = GetDescriptor(A, i); 
     float As = GetElement(tmpA, 0, featA); 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      float dif_acum = As-B.elements[k*BLOCK_SIZE + featA]; 

      accum[featA] = dif_acum[featA]*dif_acum[featA]; 
      barrier(CLK_LOCAL_MEM_FENCE); 
      for(unsigned int i = 1; i < BLOCKSIZE; i *= 2) 
      { 
       if ((featA % (2*i)) == 0) 
       accum[featA] += accum[featA + i]; 
       barrier(CLK_LOCAL_MEM_FENCE); 
      } 
      if(featA == 0) 
       C[loop] = accum[0]; 
      barrier(CLK_LOCAL_MEM_FENCE); 

      loop += 1; 
     } 
    } 
} 
+0

Je dois dire tout d'abord, que je suis très reconnaissant de la réponse de Grizzly. Je suis plutôt nouveau à OpenCL et, bien que je devais modifier le code d'exemple qu'il a donné un peu, cela m'a mené droit dans la bonne direction.Les choses importantes que j'ai remarquées (par essais et erreurs): les threads qui ne traitent pas les positions du tableau doivent être jetés; la boucle SCAN avait besoin d'un peu de peaufinage, à savoir utiliser un tampon auxiliaire pour accumuler des résultats partiels et vérifier les conditions aux limites pour trouver les termes à ajouter. Encore merci beaucoup! Je poste le code qui a fonctionné pour moi. – SebastianP

3

Merci à Grizzly, j'ai maintenant un noyau de travail. Certaines choses que j'avais besoin de modifier basé sur la réponse de Grizzly:

J'ai ajouté une instruction IF au début de la routine pour rejeter tous les threads qui ne référenceront aucune position valide dans les tableaux que j'utilise.

if(featA > BLOCK_SIZE){return;} 

Lors de la copie du premier descripteur à la mémoire locale (commune) (I.G. à Bs), l'indice doit être spécifié puisque la fonction GetElement retourne un seul élément par appel (je sautée que sur ma question).

Bs[featA] = GetElement(tmpA, 0, featA); 

Ensuite, la boucle SCAN besoin d'un peu de rangement car le tampon est remplacé après chaque itération et on ne peut pas contrôler les accès de fil les données en premier. C'est pourquoi je suis en train de 'recycler' le tampon dif_acum pour stocker des résultats partiels et ainsi éviter les incohérences tout au long de cette boucle.

dif_acum[featA] = accum[featA]; 

Il existe également un certain contrôle de limite dans la boucle SCAN pour déterminer de manière fiable les termes à additionner.

if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){ 

dernier, je pensais qu'il était logique d'inclure l'incrément variable de boucle dans la dernière instruction IF afin qu'un seul thread modifie.

if(featA == 0){ 
    C[loop] = accum[BLOCK_SIZE-1]; 
    loop += 1; 
} 

C'est tout. Je me demande toujours comment utiliser group_size pour éliminer cette définition BLOCK_SIZE et s'il y a de meilleures politiques que je peux adopter concernant l'utilisation du thread.

Ainsi, le code ressemble finalement à ceci:

__kernel void CompareDescriptors(__global float *C, DescriptorList A, DescriptorList B, int elements, __local float accum[BLOCK_SIZE], __local float Bs[BLOCK_SIZE]) 
{ 

    int gpidA = get_global_id(0); 
    int featA = get_local_id(0); 

    //global counter to store final differences 
    int loop = 0; 

    //auxiliary buffer to store temporary data 
    local float dif_acum[BLOCK_SIZE]; 

    //discard the threads that are not going to be used. 
    if(featA > BLOCK_SIZE){ 
     return; 
    } 

    //loop over all descriptors in A 
    for (int i = 0; i < A.num_elements/BLOCK_SIZE; i++){ 

     //take the gpidA-th descriptor 
     DescriptorList tmpA = GetDescriptor(A, i); 

     //copy the current descriptor to local memory 
     Bs[featA] = GetElement(tmpA, 0, featA); 

     //loop over all the descriptors in B 
     for (int k = 0; k < B.num_elements/BLOCK_SIZE; k++){ 
      //take the difference of both current descriptors 
      dif_acum[featA] = Bs[featA]-B.elements[k*BLOCK_SIZE + featA]; 

      //square the values in dif_acum 
      accum[featA] = dif_acum[featA]*dif_acum[featA]; 
      barrier(CLK_LOCAL_MEM_FENCE); 

      //copy the values of accum to keep consistency once the scan procedure starts. Mostly important for the first element. Two buffers are necesarry because the scan procedure would override values that are then further read if one buffer is being used instead. 
      dif_acum[featA] = accum[featA]; 

      //Compute the accumulated sum (a.k.a. scan) 
      for(int j = 1; j < BLOCK_SIZE; j *= 2){ 
       int next_addend = featA-(j/2); 
       if (featA >= j && next_addend >= 0 && next_addend < BLOCK_SIZE){ 
        dif_acum[featA] = accum[featA] + accum[next_addend]; 
       } 
       barrier(CLK_LOCAL_MEM_FENCE); 

       //copy As to accum 
       accum[featA] = GetElementArray(dif_acum, BLOCK_SIZE, featA); 
       barrier(CLK_LOCAL_MEM_FENCE); 
      } 

      //tell one of the threads to write the result of the scan in the array containing the results. 
      if(featA == 0){ 
       C[loop] = accum[BLOCK_SIZE-1]; 
       loop += 1; 
      } 
      barrier(CLK_LOCAL_MEM_FENCE); 

     } 
    } 
} 
+0

Je pense toujours que vous n'avez pas besoin de ces tampons locaux (accepter pour accumuler bien sûr), puisque dif_acum et Bs sont seulement accédés à la position featA, qui est l'id local du thread et donc chaque élément des tableaux est accédé par un seul fil. De plus, l'utilisation de deux tampons pour la boucle de scrutation ne devrait pas vraiment être nécessaire, puisque la cohérence est garantie par les barrières (les itérations sont séparées par des barrières et (au moins pour la boucle que j'ai posté) un fil – Grizzly

Questions connexes