Conflictos de banco de memoria compartida en CUDA: cómo se alinea la memoria con los bancos

Question

Según mi entender, la memoria compartida se divide en bancos y los accesos por múltiples hilos a un único elemento de datos dentro del mismo banco provocarán un conflicto (o emisión))

En el momento en que asignar una bastante grande array que conceptualmente representa varios pares de dos matrices:

__shared__ float A[34*N] 

Dónde N es el número de pares y los primeros 16 flotadores de un par son uno matriz y los siguientes 18 carrozas son el segundo.

El problema es que el acceso a la primera matriz no tiene conflictos, pero el acceso al segundo tiene conflictos. Estos conflictos son inevitables, sin embargo, mi pensamiento es que debido a que la segunda matriz es 18 todas las matrices futuras estarán desalineadas a los bancos y, por lo tanto, ocurrirán más conflictos de los necesarios.

¿Es esto cierto? En caso afirmativo, ¿cómo puedo evitarlo?

Cada vez que asigno la memoria compartida, ¿comienza en un banco nuevo? Así que potencialmente podría hacer

__shared__ Apair1[34] 
__shared__ Apair2[34] 
... 

¿Alguna idea?

Gracias

Answer 1

Si sus pares de matrices se almacenan de forma contigua, y si usted está accediendo a los elementos linealmente por el índice de hilo, entonces no se han compartido conflictos banco de memoria.

En otras palabras, si tiene:

A[0] <- mat1 element1 
A[1] <- mat1 element2 
A[2] <- mat1 element3 
A[15] <- mat1 element16 
A[16] <- mat2 element1 
A[17] <- mat2 element2 
A[33] <- mat2 element18 

Y acceder a este usando:

float element; 
element = A[pairindex * 34 + matindex * 16 + threadIdx.x]; 

hilos Luego adyacentes están accediendo a los elementos adyacentes en la matriz y que no tienen conflictos.

En respuesta a sus comentarios (a continuación) parece que está confundido en su comprensión. Es cierto que hay 16 bancos (en las generaciones actuales, 32 en la próxima generación, Fermi), pero las palabras consecutivas de 32 bits residen en bancos consecutivos, es decir, el espacio de direcciones se intercala entre los bancos. Esto significa que siempre que tenga un índice de matriz que puede descomponerse en x + threadIdx.x (donde x no depende de threadIdx.x, o al menos es constante en grupos de 16 subprocesos) no tendrá conflictos bancarios.

Cuando accede a las matrices más adelante en la matriz, aún tiene acceso a ellas en un bloque contiguo y, por lo tanto, no tendrá conflictos bancarios. Solo cuando empiece a acceder a elementos no adyacentes tendrá conflictos bancarios.

La reducción muestra en el SDK ilustra muy bien los conflictos bancarios mediante la construcción de una implementación ingenua a una implementación optimizada, que posiblemente vale la pena echarle un vistazo.

Answer 2

Los bancos se configuran de modo que cada 32 bits sucesivos se encuentren en el siguiente banco. Por lo tanto, si declara una matriz de flotantes de 4 bytes, cada flotante posterior en la matriz estará en el siguiente banco (módulo 16 o 32, dependiendo de su arquitectura). Asumiré que estás en capacidad de cálculo 1.x, por lo que tienes un banco de ancho 16.

Si tiene matrices de 18 y 16, las cosas pueden ser divertidas. Puede evitar conflictos de banco en la matriz de 16x16 declarándolo como

__shared__ float sixteen[16][16+1] 

que evita conflictos bancarias cuando se accede a elementos de transposición utilizando threadIdx.x (como supongo que está haciendo si usted está recibiendo los conflictos). Al acceder a elementos en, digamos, la primera fila de una matriz de 16x16, todos residirán en el primer banco. Lo que quiere hacer es tener cada uno de estos en un banco sucesivo. Padding hace esto por ti. Trate la matriz exactamente como lo haría antes, como dieciséis [fila] [columna], o de forma similar para una matriz plana, como dieciséis [fila * (16 + 1) + columna], si lo desea.

Para la caja de 18x18, al acceder en la transposición, se está moviendo a paso firme. La respuesta de nuevo es rellenar por 1.

__shared__ float eighteens[18][18+1] 

Así que ahora, cuando se accede a la transposición (decir el acceso a los elementos de la primera columna), se accederá como (1 + 18) = 16% 3, y accederá a los bancos 3, 6, 9, 12, 15, 2, 5, 8, etc., por lo que no debería tener conflictos.

El cambio de alineación particular debido a tener una matriz de tamaño 18 no es el problema, porque el punto de partida de la matriz no hace diferencia, solo es el orden en el que se accede a él. Si quieres aplanar las matrices que he propuesto anteriormente, y fusionarlas en 1, está bien, siempre y cuando tengas acceso a ellas de manera similar.