Reducciones parciales eficientes a partir de matrices de elementos, desplazamientos y longitudes de sublistas

Question

Para mi aplicación, tengo que manejar un conjunto de objetos (digamos int s) que posteriormente se divide y ordena en grupos más pequeños. Con este fin, almaceno los elementos en una única matriz continua

arr = {0,1,2,3,4,5,6,7,8,9,10,11,12,13,14...} 

y la información sobre los cubos (sublistas) viene dada por las compensaciones al primer elemento en el respectivo cubo y las longitudes de la lista secundaria.

Así, por ejemplo, dado

offsets = {0,3,8,..} 
sublist_lengths = {3,5,2,...} 

se traduciría en las siguientes divisiones:

0 1 2 || 3 4 5 6 7 || 8 9 || ... 

Lo que estoy buscando es una manera algo general y eficiente para ejecutar algoritmos, como reducciones, en los depósitos solo utilizando kernels personalizados o la biblioteca thrust. Sumando los cubos debe dar:

3 || 25 || 17 || ... 

Lo que yo he llegado con:

• opción 1: núcleos personalizados requieren un un poco de retoques, las copias en la memoria compartida, la elección correcta de bloques y tamaños de cuadrícula y una implementación propia de los algoritmos, como escanear, reducir, etc. Además, cada operación requeriría un núcleo personalizado propio. En general, es claro para mí cómo hacer esto, pero después de haber utilizado thrust en el último par de días que tengo la impresión de que podría haber una forma más inteligente

• opción 2: generar una matriz de claves de los desplazamientos ({0,0,0,1,1,1,1,1,2,2,3,...} en el ejemplo anterior) y use thrust::reduce_by_key. Aunque no me gusta la generación de listas extra.

• opción 3: Uso thrust::transform_iterator junto con thrust::counting_iterator para generar la lista de claves dado anteriormente sobre la marcha. Desafortunadamente, no puedo encontrar una implementación que no requiera incrementos de índices en la lista de compensaciones en el dispositivo y derrota el paralelismo.

¿Cuál sería la forma más sensata de implementar esto?

Answer 1

Dentro de Thrust, no puedo pensar en una solución mejor que la Opción 2. El rendimiento no será terrible, pero ciertamente no es óptimo.

Su estructura de datos tiene similitud con el formato Compressed Sparse Row (CSR) para almacenar matrices dispersas, por lo que podría utilizar técnicas desarrolladas para computar sparse matrix-vector multiplies (SpMV) para tales matrices si desea un mejor rendimiento. Tenga en cuenta que la matriz de "compensaciones" del formato CSR tiene una longitud (N + 1) para una matriz con N filas (es decir, intervalos en su caso) donde el último valor de compensación es la longitud de arr. El CSR SpMV code en Cusp es un poco intrincado, pero sirve como un buen punto de partida para su kernel. Simplemente elimine cualquier referencia a Aj o x del código y pase offsets y arr en los argumentos Ap y Av, respectivamente.

Answer 2

No mencionó qué tan grandes son las cubetas. Si los cubos son lo suficientemente grandes, tal vez pueda salirse con la suya copiando los desplazamientos y sublista_lengths en el host, iterando sobre ellos y haciendo una llamada a Thrust por separado para cada segmento. Fermi puede tener 16 núcleos en vuelo al mismo tiempo, por lo que en esa arquitectura es posible que pueda manejar cubos más pequeños y aún obtener una buena utilización.