OpenCL - suma incremental durante el cálculo

Question

Soy absolutamente novato en programación OpenCL. Para mi aplicación (simulaton molecular) Escribí un kernel para calcular el potencial intermolecular del líquido de lennard-jones. En este núcleo que necesito para calcular el valor acumulativo del potencial de todas las partículas con una:

__kernel void Molsim(__global const float* inmatrix, __global float* fi, const int c, const float r1, const float r2, const float r3, const float rc, const float epsilon, const float sigma, const float h1, const float h23) 
{ 
    float fi0; 
    float fi1; 
    float d; 

    unsigned int i = get_global_id(0); //number of particles (typically 2000) 

    if(c!=i) { 
     // potential before particle movement 
     d=sqrt(pow((0.5*h1-fabs(0.5*h1-fabs(inmatrix[c*3]-inmatrix[i*3]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(inmatrix[c*3+1]-inmatrix[i*3+1]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(inmatrix[c*3+2]-inmatrix[i*3+2]))),2.0)); 
     if(d<rc) { 
     fi0=4.0*epsilon*(pow(sigma/d,12.0)-pow(sigma/d,6.0)); 
     } 
     else { 
     fi0=0; 
     } 
     // potential after particle movement 
     d=sqrt(pow((0.5*h1-fabs(0.5*h1-fabs(r1-inmatrix[i*3]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(r2-inmatrix[i*3+1]))),2.0)+pow((0.5*h23-fabs(0.5*h23-fabs(r3-inmatrix[i*3+2]))),2.0)); 
     if(d<rc) { 
     fi1=4.0*epsilon*(pow(sigma/d,12.0)-pow(sigma/d,6.0)); 
     } 
     else { 
      fi1=0; 
     } 
     // cumulative difference of potentials 
     // fi[0]+=fi1-fi0; changed to full size vector 
     fi[get_global_id(0)]=fi1-fi0; 
     } 
}   

Mi problema está en la línea de: fi [0] + = fi1-fi0 ;. En el vector de un elemento fi [0] hay resultados incorrectos. Leí algo sobre la reducción de suma, pero no sé cómo hacerlo durante el cálculo.

¿Existe alguna solución simple a mi problema?

Aviso: He intentado añadir al lado del núcleo para la suma de las componentes del vector (véase el código de abajo), pero no hubo una desaceleración aún mayor que cuando suma vectorial utilizando la CPU.

__kernel void Arrsum(__global const float* inmatrix, __global float* outsum, const int inmatrixsize, __local float* resultScratch) 
{ 
     // načtení indexu 
     int gid = get_global_id(0); 
     int wid = get_local_id(0); 
     int wsize = get_local_size(0); 
     int grid = get_group_id(0); 
     int grcount = get_num_groups(0); 

     int i; 
     int workAmount = inmatrixsize/grcount; 
     int startOffest = workAmount * grid + wid; 
     int maxOffest = workAmount * (grid + 1); 
     if(maxOffest > inmatrixsize){ 
     maxOffest = inmatrixsize; 
    } 

    resultScratch[wid] = 0.0; 
    for(i=startOffest;i<maxOffest;i+=wsize){ 
      resultScratch[wid] += inmatrix[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 

    if(gid == 0){ 
      for(i=1;i<wsize;i++){ 
        resultScratch[0] += resultScratch[i]; 
      } 
      outsum[grid] = resultScratch[0]; 
    } 
} 

Answer 1

Creo que necesita la función atómica atomic_add para fi [0] + = fi1-fi0;

Advertencia: el uso de una función atómica reduce el rendimiento.

Aquí, dos ejemplos con la función atómica de incremento.

Ejemplo sin función atómica y 2 workitems:

__kernel void inc(global int * num){ 
    num[0]++; //num[0] = 0 
} 

1. de elemento de trabajo 1 lee num [0]: 0
2. de elemento de trabajo 2 lee num [0]: 0
3. Trabajo El artículo 1 incrementa num [0]: 0 + 1
4. El ítem de trabajo 2 incrementa num [0]: 0 + 1
5. El ítem de trabajo 1 escribe num [0] : Num [0] = 1
6. de elemento de trabajo 2 escribe num [0]: num [0] = 1

Resultado: num [0] = 1

Ejemplo con función atómica y 2 workitems:

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable 

__kernel void inc(global int * num){ 
    atom_inc(&num[0]); 
} 

1. de elemento de trabajo 1 lee num [0]: 0
2. de elemento de trabajo incrementos de 1 num [0]: 0 + 1
3. de elemento de trabajo 1 escribe num [0]: num [0] = 1
4. de elemento de trabajo 2 lee num [0]: 1
5. de elemento de trabajo incrementos de 2 num [0]: 1 + 1
6. Trabajo artículo 2 escribe num [0]: num [0] = 2

Resultado: num [0] = 2

Answer 2

complemento Atómica es una solución, pero se podía conseguir los problemas de rendimiento, ya que la parte atómica serializará tus artículos de trabajo

Creo que la mejor solución es, para cada elemento de trabajo, a escribir en su propia variable, como:

fi [get_global_id (0)] + = fi1-fi0;

Luego puede transferir la matriz a la CPU y sumar todos los elementos, o puede hacerlo en la GPU con un algoritmo para hacerlo en paralelo.

Answer 3

Todos los hilos son ejecutados por "grupos". Puede determinar el id. Del subproceso en el grupo utilizando la función get_local_id (dim). Los hilos dentro de cada grupo pueden usar memoria compartida (que se llama "memoria local" en OpenCL) y sincronizar su ejecución, pero los hilos en diferentes grupos no pueden comunicarse directamente.

Por lo tanto, la solución típica para la reducción es el siguiente:

1. Añadir part_sum temporal matrices (global) y tmp_reduce (local) para los argumentos del kernel:

   __kernel void Molsim(..., __global float *part_sum, __local float *tmp_reduce) 
   

2. Asignar serie de flotadores con un tamaño igual al número de grupos (= global_size/local_size) de su kernel y establecer el parámetro part_sum.

3. Conjunto de parámetros tmp_reduce con su kernel "de tamaño local de" x size_of (float) y NULL:

   clSetKernelArg(kernel,<par number>,sizeof(float)*<local_size>,NULL); 
   

4. En el núcleo, reemplace el código con lo siguiente:

   int loc_id=get_local_id(0); 
   
   ... 
   
   //  fi[0]+=fi1-fi0; 
        tmp_reduce[loc_id]=fi1-fi0; 
        } 
       barrier(CLK_LOCAL_MEM_FENCE); 
       if(loc_id==0) { 
       int i; 
       float s=tmp_reduce[0]; 
       for(i=1;i<get_local_size(0);i++) 
        s+=tmp_reduce[i]; 
       part_sum[get_group_id(0)]=s; 
       } 
   } 
   

5. Después finalizando la ejecución del kernel, simplemente sume en el host los contenidos de part_sum [array], que es mucho más pequeño, que global_size.

Esto no es totalmente "reducción paralela", porque se puede resumir tmp_reduce matriz en operaciones paralelas utilizando log2 (local_size) utilizando algoritmos más complicados, pero esto debe ser mucho más rápido que las operaciones atómicas.

Además, consulte http://developer.amd.com/Resources/documentation/articles/pages/OpenCL-Optimization-Case-Study-Simple-Reductions_2.aspx para ver los métodos de reducción paralelos de Beter.