Acceso a la memory intermedia de la GPU por índice

Nota: Mi pregunta se refiere a Apple's Metal API, pero creo que el concepto es lo suficientemente general como para traducir a otros frameworks de GPU.

Mi objective: agregar un vector b fila 1 x N a cada fila en una matriz M x N A

Mi kernel, networkingucido a la parte con la que estoy teniendo problemas:

 kernel void vmadd(const device float* A [[ buffer(0) ]], const device float* b [[ buffer(1) ]], device float* C [[ buffer(2) ]], constant ushort& aWidth [[ buffer(3) ]], ushort2 gid [[ thread_position_in_grid ]]) { int idx = gid.y * aWidth + gid.x; // Compute absolute index in C C[idx] = A[idx] + b[gid.x]; } 

Suposition: entiendo que gid es la position de un solo elemento en C : gid.x es la columna, gid.y es la fila. Si esto no es cierto, alguien me corrige.

Ahora si completo A con 8 x 8 ceros:

 A = 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 

y b así:

 b = 1 2 3 4 5 6 7 8 

luego de la ejecución, C debería ser una matriz de 8 x 8 donde cada fila es 1 2 3 4 5 6 7 8 .

En cambio, obtengo esto:

 C = 1 2 3 4 5 6 7 8 1 2 3 4 5 6 7 8 1 2 3 4 5 6 7 8 0 0 0 0 0 0 0 0 1 2 3 4 5 6 7 8 1 2 3 4 5 6 7 8 1 2 3 4 5 6 7 8 0 0 0 0 0 0 0 0 

Estoy evaluando grupos de subprocesss de acuerdo con la recomendación de Apple aquí :

 let w = computePipeline.threadExecutionWidth let h = computePipeline.maxTotalThreadsPerThreadgroup / w let threadsPerThreadgroup = MTLSizeMake(w, h, 1) let threadgroupsPerGrid = MTLSize(width: (cWidth + w - 1) / w, height: (cHeight + h - 1) / h, depth: 1) 

Lo que en mi máquina produce hilos (64, 16, 1) por grupo y (1, 1, 1) grupos de hilos por cuadrícula.

Sin embargo, si pongo manualmente threadsPerThreadgroup en (1, 1, 1) y threadgroupsPerGrid en (8, 8, 1) , obtengo el resultado correcto en C

La pregunta:

Estoy bastante seguro de que mi problema tiene que ver con el tamaño del subgrupo y la manera en que se gestionan los búferes, pero soy bastante nuevo en la progtwigción de la GPU, así que no lo entiendo completamente.

¿Por qué disminuir el tamaño de los grupos de hilos produce el resultado correcto? Incluso más en general, ¿por qué el dimensionamiento de grupos de hilos tiene algún efecto en este cálculo en absoluto?

Me parece que si gid siempre corresponde a un índice en C , y le pido a b[gid.x] , ese valor debería estar disponible. Entiendo que los datos generalmente se dividen para caber en cachings de grupos de subprocesss; si ese es el caso aquí, ¿qué reglas sigue y cómo puedo explicarlo?

Mi entendimiento es que gid es la position de un solo elemento en C : gid.x es la columna, gid.y es la fila. Si esto no es cierto, alguien me corrige.

Esto no es completamente cierto. gid es la position dentro de la cuadrícula .

Debido a que la cuadrícula es, como sucede, 64×16, la function de cálculo será llamada para las posiciones que caen fuera de sus matrices 8×8 ( A y C ) y su vector de 8 elementos ( b ). Cuando eso sucede, las lecturas de A pueden acceder a la fila incorrecta o incluso más allá del final de A De manera similar, las lecturas de b leerán más allá de su final.

Por ejemplo, considere cuando gid es (8, 0). idx será 8. idx A[8] , que en realidad está en (0, 1). Leerás b[8] , que es pasado el final. Eso es técnicamente indefinido, pero es muy probable que sea 0 para un búfer de esa longitud relativamente corta. Escribirá en C[8] que también está en (0, 1). Esto está sucediendo aproximadamente simultáneamente como la invocación de function que se supone que debe escribir en (0, 1) y hay una raza que prevalece.

Su function debería, cerca del comienzo, probar si gid está fuera de límites y, de ser así, regresar temprano:

 if (any(gid > aWidth)) return; 

(Esto supone que A y C siempre serán cuadrados, por lo que el ancho y el alto se pueden comprobar con el único valor).

Puede intentar ajustar los cálculos de threadsPerThreadgroup y threadgroupsPerGrid para que la cuadrícula sea exactamente del tamaño de sus matrices, pero puede ser tedioso hacerlo correctamente en todos los casos. Dicho esto, ciertamente puede evitar que threadsPerThreadgroup sea ​​demasiado grande:

 let w = min(computePipeline.threadExecutionWidth, cWidth) let h = min(computePipeline.maxTotalThreadsPerThreadgroup / w, cHeight) 

Pero aún necesitará la verificación en la function de cálculo porque la cuadrícula total aún puede ser demasiado grande. Por ejemplo, supongamos que computePipeline.threadExecutionWidth es al less 8 y computePipeline.maxTotalThreadsPerThreadgroup es 60. Bien, w será 8 pero h será 7. Entonces, threadgroupsPerGrid será (1, 2, 1) y el tamaño total de la cuadrícula será 8x14x1 , que es nuevamente más grande que sus matrices.