2010-09-22 10 views
5

Estoy calculando la distancia euclidiana entre puntos n-dimensionales usando OpenCL. Obtengo dos listas de puntos n-dimensionales y debería devolver una matriz que contiene solo las distancias desde cada punto en la primera tabla hasta cada punto en la segunda tabla.Suma acumulativa acumulativa usando OpenCL

Mi enfoque es hacer el bucle doble regular (por cada punto en la Tabla 1 {por cada punto en la Tabla2 {...}} y luego hacer el cálculo para cada par de puntos en paralelo.

El euclidiana la distancia se divide en 3 partes: 1. tome la diferencia entre cada dimensión en los puntos 2. cuadre esa diferencia (aún para cada dimensión) 3. sume todos los valores obtenidos en 2. 4. Tome la raíz cuadrada del valor obtenido en 3. (este paso se ha omitido en este ejemplo)

Todo funciona como un amuleto hasta que intente acumular la suma de todas las diferencias (es decir, ejecutar el paso 3. del procedimiento descrito anteriormente, línea 49 del código a continuación).

Como datos de prueba estoy usando DescriptorLists con 2 puntos cada uno: 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)

Así que el vector resultante debería tener los valores: 128, 2064512, 2130048, 128 Ahora mismo estoy obteniendo números aleatorios que varían con cada ejecución.

Agradezco cualquier ayuda o pistas sobre lo que estoy haciendo mal. Con suerte todo está claro el escenario estoy trabajando en

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

Respuesta

7

Su problema radica en estas líneas de código:.

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

Todos los temas en su grupo de trabajo (bueno, en realidad todos sus hilos, pero vayamos a eso más adelante) estamos tratando de modificar esta posición de la matriz al mismo tiempo sin ninguna sincronización de ningún tipo. Hay varios factores que hacen que esto realmente problemático:

  1. El grupo de trabajo no se garantiza que funcione completamente en paralelo, lo que significa que para algunos hilos C [LOOP] = 0 puede ser llamado después de que otros temas que ya han ejecutado la siguiente línea
  2. Aquellos que ejecutan en paralelo, todos leen el mismo valor de C [bucle], lo modifican con su incremento e intentan volver a escribir en la misma dirección. No estoy completamente seguro de cuál es el resultado de esa reescritura (creo que uno de los hilos tiene éxito al escribir de nuevo, mientras que los otros fallan, pero no estoy del todo seguro), pero está mal de cualquier manera.

Ahora vamos a solucionar este problema: pesar de que podría ser capaz de conseguir que esto funcione en la memoria global usando atómicas, no va a ser rápido, por lo que permite acumular en la memoria local:

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 Por supuesto, puede reutilizar otros almacenamientos intermedios locales para esto, pero creo que el punto es claro (por cierto: ¿Está seguro de que dif_acum se creará en la memoria local, porque creo que leí en alguna parte que esto no se colocaría en la memoria local, haría que la precarga de A en la memoria local fuera algo inútil).

Algunos otros puntos sobre este código:

  1. Su código es parece estar orientado a usar sólo en grupo de trabajo (que no está utilizando ya sea Identificación del groupid ni global para ver qué elementos para trabajar en), por un rendimiento óptimo que quizás desee utilizar más que eso.
  2. podría ser preferance personal, pero me parece mejor utilizar get_local_size(0) para la workgroupsize que utilizar un Definir (puesto que podría cambiar en el código de host sin darse cuenta de que debería haber cambiado su código OpenCL a)
  3. Las barreras en su código son todas innecesarias, ya que ningún hilo accede a un elemento en la memoria local que está escrito por otro hilo. Por lo tanto, no necesita usar memoria local para esto.

Teniendo en cuenta la última bala que podría simplemente hacer:

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

Esto haría que el código (sin considerar los dos primeros puntos):

__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

Tengo que decir, en primer lugar, que estoy muy agradecido con la respuesta de Grizzly. Soy bastante nuevo en OpenCL y, aunque necesitaba modificar el código de ejemplo que me dio un poco, me condujo directamente a la dirección correcta.Cosas importantes que noté (por prueba y error): hilos que no abordan las posiciones de la matriz necesitaban ser descartados; el bucle SCAN necesitaba un pequeño ajuste, a saber, el uso de un búfer auxiliar para acumular resultados parciales y verificar las condiciones de contorno para encontrar los términos que se agregarían. ¡De nuevo, muchas gracias! Estoy publicando el código que funcionó para mí. – SebastianP

3

Gracias a grisáceo, que tengo ahora un kernel en funcionamiento Algunas cosas que necesitaba modificar basadas en la respuesta de Grizzly:

Agregué una instrucción IF al comienzo de la rutina para descartar todos los hilos que no hagan referencia a ninguna posición válida en las matrices que estoy usando.

if(featA > BLOCK_SIZE){return;} 

Al copiar el primer descriptor de memoria local (compartido) (i.g. a B), el índice se tiene que especificar ya que la función GetElement devuelve sólo un elemento por llamada (I omiten que en mi pregunta).

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

Entonces, el bucle SCAN necesita un poco de ajuste porque el búfer se sobrescribe después de cada iteración y no se puede controlar el acceso que se enroscan los datos primero. Es por eso que estoy 'reciclando' el buffer de dif_acum para almacenar resultados parciales y de esa manera, prevenir inconsistencias a lo largo de ese ciclo.

dif_acum[featA] = accum[featA]; 

Hay también algunos de control de límites en el bucle SCAN para determinar de forma fiable los términos que se añaden juntos.

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

pasado, pensé que tenía sentido para incluir el mínimo de la variable de bucle dentro de la última instrucción SI de modo que sólo un hilo modifica.

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

Eso es todo. Todavía me pregunto cómo puedo usar group_size para eliminar esa definición BLOCK_SIZE y si hay mejores políticas que puedo adoptar con respecto al uso de subprocesos.

Así que el código es el fin de esta manera:

__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

Sigo pensando que no necesita esos búferes locales (acepte para acum por supuesto), ya que a ambos dif_acum y Bs solo se accede en la posición featA, que es la identificación local del hilo y, por lo tanto, se accede a cada elemento de los arrays solo un hilo. Además, el uso de dos búferes para el bucle de exploración no debería ser realmente necesario, ya que la consistencia está garantizada por las barreras (las iteraciones están separadas por barreras y (al menos para el bucle que publiqué) en cada iteración, cada elemento solo es accedido por un hilo – Grizzly

Cuestiones relacionadas