Primer lugar, permítanme decir que soy plenamente consciente de que mi pregunta ya ha sido pedido: Bloque reducción en CUDA sin Embargo, como espero hacer claro, mi pregunta es un seguimiento y tengo necesidades particulares que hacen que la solución encontrada por que OP sea adecuado.

Así, me explico. En mi actual código, tengo un kernel Cuda en cada iteración de un bucle while para hacer algunos cálculos sobre los valores de una matriz. Como un ejemplo, piensa que es como la siguiente:

int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
    __global__ void calcKernel(int* idata, int* odata)
    {
        int i = blockIdx.x*blockDim.x + threadIdx.x;
        if (i < n)
        {
            odata[i] = (idata[i] + 2) * 5;
        }
    }

    iteration++;
}

Sin embargo, la siguiente que tengo que cumplir aparentemente difícil tarea de la GPU. En cada iteración del bucle while que se llama el kernel, tengo que suma todos los valores generados dentro de odata y guardar el resultado en un intarray llamado result, en una posición dentro de la matriz que corresponde a la iteración actual. Tiene que ser logrado en el interior del núcleo o al menos todavía en la GPU debido a que limita el rendimiento, sólo puedo recuperar el result matriz en la final después de todas las iteraciones se ha completado.

Un mal ingenuo intento de woud el siguiente aspecto:

int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
    __global__ void calcKernel(int* idata, int* odata, int* result)
    {
        int i = blockIdx.x*blockDim.x + threadIdx.x;
        if (i < n)
        {
            odata[i] = (idata[i] + 2) * 5;
        }
    }

    result[iteration] = 0;
    for(int j=0; j < max_iterations; j++)
    {
        result[iteration] += odata[j];            
    }

    iteration++;
}

Por supuesto, el código anterior no funciona debido a la GPU distribuir el código a través de los subprocesos. Con el fin de lear en la manera correcta de hacerlo, he estado leyendo otras preguntas aquí en el sitio sobre la matriz de reducción de uso de CUDA. En particular, he encontrado una mención a un muy buen NVIDIA pdf sobre tal tema, que también se discute en la antigua ASÍ que la pregunta que mencioné al principio: http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

Sin embargo, aunque estoy totalmente de entender los pasos de el código que se describen en dicha diapositivas, así como, en general, optimizaciones, no entiendo cómo ese enfoque se puede resumir-reducir una matriz a un número si el código que realmente ouptus una gama completa (y uno de claro dimensiones). Podría alguien por favor, arrojar algo de luz sobre esto y me muestran un ejemplo de la forma en que funcione (es decir, cómo conseguir el número de la matriz de salida)?

Ahora, volviendo a la pregunta que he mencionado al principio (Bloque reducción en CUDA). Tenga en cuenta que su aceptada respuesta solo sugiere leer el pdf que he enlazado más arriba – que no no hablar acerca de qué hacer con la matriz de salida generadas por el código. En los comentarios, el OP no menciona que él/ella fue capaz de finishi el trabajo haciendo un resumen de la matriz de salida en la CPU, que es algo que no puedo hacer, ya que significaría la descarga de la matriz de salida de cada iteración de mi bucle while. Por último, la tercera respuesta en ese enlace sugiere el uso de una biblioteca para realizar esto, pero estoy interesado en aprender el modo nativo de hacerlo.

Por otra parte, que sería también muy interesado en cualquier otro proposiciones acerca de cómo implementar lo que estoy descrito anteriormente.

  • Habría que considerar el uso de varios núcleos, en lugar de uno solo?
  • Seguro, siempre que no me transferencia de/para la memoria de la GPU, llamar a más de un núcleo no será un problema per se.
  • Yo recomendaría a continuación, utilice 2 núcleos, 1 para los primeros cálculos, y otro kernel, solo para agregar, y colocar el bucle, en el programa de Acogida. Necesito más información de cómo el programa funciona, pero con el código que nos muestran, que podría ser una manera fácil de resolver su problema.
  • Ahora estoy probando solo eso, sino que estoy luchando con el último paso. Exactamente lo mismo que pidió originalmente en la pregunta que he enlazado (pero que no fue resuelto en el lado de la CPU). Es, estoy usando un kernel ahora para hacer la reducción, pero termino con un array, no un único número, y quiero reducir aún más este final de la matriz de un número único dentro de la GPU.
  • En el ejemplo que les he dado, el resultado en el primer elemento de la matriz, si lo quieres en una sola variable (lugar de la memoria) enviar un puntero en el núcleo, y asignar ese valor a la ubicación donde desea que el valor, de esa manera, cuando se desea leer, desde el Host, o el Dispositivo, usted puede darle a ese puntero.
  • Se consideraría la posibilidad de ‘Empuje’ de la API para hacer la suma para usted? Se separa la funcionalidad y fácil en su uso. Una llamada a la API en la GPU no ser que mucho tiempo desde mi propia experiencia. Decenas de microsegundos debe hacer el trabajo.

InformationsquelleAutor AndrewSteer | 2017-03-01

2 Comentarios

  1. 4

    Que ya se ha encontrado la canónica información sobre el bloque de reducciones paralelas, así que no voy a repetir eso. Si usted no desea escribir un montón de código nuevo para hacer esto, sugiero que se busca en el CUB de la biblioteca block_reduce aplicación, que proporciona una óptima bloque sabio operación de reducción con la adición de alrededor de 4 líneas de código al kernel existente.

    Sobre la verdadera pregunta aquí, usted puede hacer lo que usted desea si usted hacer algo como esto:

    __global__ void kernel(....., int* iter_result, int iter_num) {
    
        //Your calculations first so that each thread holds its result
    
        //Block wise reduction so that one thread in each block holds sum of thread results
    
        //The one thread holding the adds the block result to the global iteration result
        if (threadIdx.x == 0)
            atomicAdd(iter_result + iter_num, block_ressult);
    }

    La clave aquí es que un atómica función se utiliza de forma segura actualizar el kernel ejecutar resultado con los resultados de un bloque dado, sin un recuerdo de la carrera. Usted absolutamente debe inicializar iter_result antes de ejecutar el kernel, de lo contrario el código no funciona, sino que es el núcleo básico de patrones de diseño.

    • Muy perspicaz. Yo sólo era capaz de aplicar correctamente y que funciona. El único inconveniente en mi caso es que estoy trabajando con caracteres y atomicAdd no funciona con caracteres, así que he tenido que convertir las partes relevantes, perdiendo un poco de memoria y un poco de rendimiento. Aún así, en un puro comparación «atomicAdd» frente a «ahorrar bloque de resultados en la matriz de salida», estoy sorprendido por el poco rendimiento se perdió por el uso de atomicAdd.
  2. 2

    Si agrega 2 números contiguos, y guardar el resultado, en cualquiera de las ranuras donde guardar los números, usted sólo tiene que ejecutar varias veces el mismo kernel, para mantener la reducción en 2 potencia de la matriz de sumas de dinero, como en este ejemplo:

    De la matriz de suma de los valores:

    12345678910]

    Primera ejecución n/2 hilos, suma contiguos de los elementos de la matriz, y la almacenan en la «izquierda» de cada uno, la matriz se verá ahora como:

    3,27,411,615,819,10]

    Ejecutar el mismo núcleo del sistema, ejecute n/4 hilos, ahora agregar cada uno de los 2 elementos, y de la tienda a la izquierda, la mayoría de los elemento de la matriz de ahora se verá así:

    10,2,7,426,6,15,819,10]

    Ejecutar el mismo núcleo del sistema, ejecute n/8 hilos, ahora agregar cada uno de los 4 elementos, y de la tienda en la parte izquierda de la mayoría de los elemento de la matriz, para obtener:

    36,2,7,4,26,6,15,819,10]

    Ejecutar una última vez, un solo hilo para agregar cada uno de los 8 elementos, y de la tienda en la parte izquierda de la mayoría de los elemento de la matriz, para obtener:

    [55,2,7,4,26,6,15,8,19,10]

    De esta manera, usted sólo tiene que ejecutar el kernel con algunos hilos como parámetros, para obtener la redux al final, en el primer elemento (55) mirar a los «puntos» (·) para ver que los elementos de la matriz se «activa» a la suma de ellos, de cada ejecución.

    • Gracias por tu respuesta. Aunque entiendo la lógica, no veo cómo implementar eso. Me podría dar un ejemplo insignificante? Parece que estás hablando exactamente de lo mismo que se propuso en el pdfhe vinculado en la pregunta.

Dejar respuesta

Please enter your comment!
Please enter your name here