reducción de openCL, y paso 2d array

Question

Aquí está el bucle que quiero convertir a openCL.

for(n=0; n < LargeNumber; ++n) {  
    for (n2=0; n2< SmallNumber; ++n2) { 
     A[n]+=B[n2][n]; 
    }               
    Re+=A[n];  
} 

Y esto es lo que tengo hasta ahora, aunque, sé que no es correcto y me faltan algunas cosas.

__kernel void openCL_Kernel(__global int *A, 
         __global int **B, 
         __global int *C, 
         __global _int64 Re, 
            int D) 
{ 

int i=get_global_id(0); 
int ii=get_global_id(1); 

A[i]+=B[ii][i]; 

//barrier(..); ? 

Re+=A[i]; 

} 

Soy un principiante completo para este tipo de cosas. Antes que nada, sé que no puedo pasar un doble puntero global a un kernel openCL. Si puede, espere unos días más o menos antes de publicar la solución, quiero resolver esto por mí mismo, pero si puede ayudarme a orientarme en la dirección correcta, le estaría agradecido.

Answer 1

Acerca de su problema con pasar puntos dobles: ese tipo de problema normalmente se resuelve copiando toda la matriz (o lo que sea que esté trabajando) en un bloque continuo de memoria y, si los bloques tienen diferentes longitudes pasando otra matriz, que contiene los desplazamientos para las filas individuales (por lo que su acceso se vería como B[index[ii]+i]).

Ahora para su reducción a Re: ya que no mencionó en qué tipo de dispositivo está trabajando voy a asumir su GPU. En ese caso, evitaría hacer la reducción en el mismo kernel, ya que será lento como en el modo en que lo publicaste (tendrías que serializar el acceso al Re en miles de hilos (y también al acceso al A[i]). En cambio, escribiría kernel de búsqueda, que suma todos B[*][i] en A[i] y puse la reducción de A en Re en otro kernel y lo hago en varios pasos, es decir, utiliza un kernel de reducción que opera en el elemento n y los reduce a algo como n/16 (o cualquier otro número). Luego llama iterativamente ese kernel hasta que esté abajo a un elemento, que es el resultado (estoy haciendo esta descripción intencionalmente vaga, ya que dijo que quería pensar en usted mismo).

Como nota al margen: ¿Se da cuenta de que el código original no tiene exactamente un buen patrón de acceso a la memoria? Suponiendo que B es relativamente grande (y mucho más grande que A debido a la segunda dimensión) tener el ciclo interno iterar sobre el índice externo va a crear una gran cantidad de memorias caché. Esto es aún peor cuando al portar a la GPU, que es muy sensible acerca de acceso a la memoria coherente

Así reordenación de ello como esto puede aumentar enormemente el rendimiento:

for (n2=0; n2< SmallNumber; ++n2) 
    for(n=0; n < LargeNumber; ++n)  
    A[n]+=B[n2][n]; 
for(n=0; n < LargeNumber; ++n)             
    Re+=A[n];  

Esto es particularmente cierto si usted tiene un compilador que es bueno en autovectorización, ya que podría ser capaz de vectorizar esa construcción, pero es muy poco probable que pueda hacerlo para el código original (y si no puede probar que A y B[n2] no pueden referirse a la misma memoria que puede convertir el código original en esto).