Uso de la variable del dispositivo por múltiples hilos en CUDA

Question

Estoy jugando con Cuda.

Por el momento tengo un problema. Estoy probando una gran matriz para respuestas particulares, y cuando obtengo la respuesta, tengo que copiar los datos en otra matriz.

Por ejemplo, mi matriz de prueba de 5 elementos se ve así:
] [] [v1] [] [] [v2

El resultado debe verse así:
V1] [V2

El problema es ¿cómo calculo la dirección de la segunda matriz para almacenar el resultado? Todos los elementos de la primera matriz se verifican en paralelo.

Estoy pensando en declarar una variable de dispositivo int addr = 0. Cada vez que encuentro una respuesta, incrementaré el addr. Pero no estoy seguro de eso porque significa que addr Se puede acceder mediante múltiples hilos al mismo tiempo. ¿Eso causará problemas? ¿O el hilo esperará hasta que otro hilo termine con esa variable?

Answer 1

No es tan trivial como parece. Acabo de terminar para implementar uno y puedo decir lo que necesitas leer el Artículo de escaneo de GPU Gems 3 en particular capítulo 39.3.1 Compactación de flujo.

Para implementar su propio comienzo desde el ejemplo LargearRayScan en el SDK, eso le dará solo el prescan. Suponiendo que tenga la matriz de selección en la memoria del dispositivo (una matriz de 1 y 0, lo que significa 1- Seleccione 0 - Descarte), dev_selection_array a dev_elements_array elementos a seleccionar un dev_prescan_array y un dev_result_array todo tamaño norte entonces tu haces

prescan(dev_prescan_array,dev_selection_array, N);
scatter(dev_result_array, dev_prescan_array,
         dev_selection_array, dev_elements_array, N);

donde esta la dispersión está

 __global__ void scatter_kernel( T*dev_result_array, 
                   const T* dev_prescan_array, 
                   const T* dev_selection_array,
                   const T* dev_elements_array, std::size_t size){

unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= size) return;
if (dev_selection_array[idx] == 1){
    dev_result_array[dev_prescan_array[idx]] = dev_elements_array[idx];
}
}

Para otra buena aplicación del prescan, consulte el documento BLE93

¡Que te diviertas!

Answer 2

Estás hablando de la compactación clásica de la transmisión. Generalmente recomendaría mirar Empuje o CUDPP (Esos enlaces van a la documentación de compactación). Ambos son de código abierto, si desea rodar el suyo, también sugeriría ver la muestra SDK 'Scan'.