Conflicto de banco de memoria compartida GPU

Question

Estoy tratando de comprender cómo se producen los conflictos bancarios.
si tengo una matriz de 256 en la memoria global y tengo 256 hilos en un solo bloque, y quiero copiar la matriz en la memoria compartida. por lo tanto, cada hilo copia un elemento.

shared_a[threadIdx.x]=global_a[threadIdx.x] 

¿esta acción simple resulta en un conflicto bancario?

Supongamos ahora que el tamaño de la matriz es mayor que el número de hilos, por lo que ahora estoy utilizando esto para copiar la memoria global de la memoria compartida:

tid = threadIdx.x; 
for(int i=0;tid+i<N;i+=blockDim.x) 
    shared_a[tid+i]=global_a[tid+i]; 

hace el resultado código anterior en una conflicto bancario?

Answer 1

La mejor manera de comprobar esto sería perfilar su código con el "Compute Visual Profiler"; esto viene con CUDA Toolkit. También hay una gran sección en GPU Gems 3 sobre esto - "39.2.3 Evitando Conflictos de Banco".

"Cuando varios hilos en el mismo acceso a la deformación del mismo banco, un conflicto banco se produce a menos que todos los hilos del acceso a la deformación de la misma dirección dentro de la misma palabra de 32 bits" - Lo primero que hay 16 bancos de memoria cada uno 4bytes de ancho Así que, esencialmente, si usted tiene cualquier tema en un medio urdimbre lectura de la memoria de las mismas 4bytes en un banco de memoria compartida, vas a tener conflictos bancarias y serialización etc.

OK para que el primer ejemplo:

en primer lugar permite asumen sus matrices son digamos por ejemplo del tipo int (una palabra de 32 bits). Su código guarda estas entradas en la memoria compartida, en cualquier medio warp el hilo Kth está guardando en el banco de memoria Kth. Entonces, por ejemplo, el subproceso 0 de la primera mitad de urdimbre se guardará en shared_a[0] que está en el primer banco de memoria, el subproceso 1 se guardará en shared_a[1], cada mitad de urdimbre tiene 16 subprocesos estos mapas en los 16 bancos de 4 bytes. En la siguiente mitad warp, el primer hilo guardará ahora su valor en shared_a [16] que está en el primer banco de memoria . Entonces, si usa una palabra de 4 bytes como int, float, etc., su primer ejemplo no generará un conflicto bancario. Si usa una palabra de 1 byte como char, en la primera mitad los hilos de urdimbre 0, 1, 2 y 3 guardarán todos sus valores en el primer banco de memoria compartida que causará un conflicto bancario.

Segundo ejemplo:

Una vez más todo esto va a depender del tamaño de la palabra que está utilizando, pero para el ejemplo voy a usar una palabra 4byte. Así que buscando en la primera mitad de la urdimbre:

Número de hilos = 32

N = 64

Tema 0: escribirá a 0, 31, 63 Tema 1: escribirá a 1, 32

Todos los subprocesos de la mitad de la disformidad se ejecutan simultáneamente, por lo que las escrituras en la memoria compartida no deben causar conflictos bancarios. Sin embargo, tendré que revisar esta otra vez.

Espero que esto ayude, perdón por la gran respuesta!

Answer 2

En ambos casos, los subprocesos acceden a la memoria compartida con dirección consecutiva. Depende del tamaño del elemento de la memoria compartida, pero el acceso consecutivo a la memoria compartida por una urdimbre de hilos no da como resultado un conflicto bancario para tamaños de elementos "pequeños".

Perfilando this code con NVIDIA Visual Profiler muestra que para un tamaño de elemento menor que 32 y un múltiplo de 4 (4, 8, 12, ..., 28), el acceso consecutivo a la memoria compartida no genera un conflicto bancario . El tamaño del elemento de 32, sin embargo, da como resultado un conflicto bancario.

---

respuesta por Ljdawson contiene alguna información obsoleta:

... Si utiliza una palabra de 1 byte tales como char, en la primera mitad hilos de urdimbre 0, 1, 2 y 3, todos se guarde sus valores en el primer banco de memoria compartida que causará un conflicto bancario.

Esto puede ser cierto para las GPU de edad, pero para los últimos GPUs con CC> = 2.x, que no causan conflictos de banco, de manera efectiva debido al mecanismo de difusión (link). La siguiente cita es de CUDA C PROGRAMMING GUIDE (v8.0.61) G3.3. Shared Memory.

Una solicitud de memoria compartida para un warp no genera un conflicto bancario entre dos hilos que acceden a cualquier dirección dentro de la misma palabra de 32 bits (aunque las dos direcciones caen en el mismo banco): En ese caso, para los accesos de lectura, la palabra se transmite a los hilos solicitantes (se pueden transmitir múltiples palabras en una sola transacción) y para los accesos de escritura, cada dirección está escrita solo por uno de los hilos (qué hilo realiza la escritura no está definida).

Esto significa, en particular, que no existen conflictos de banco si se accede a una matriz de char como sigue, por ejemplo:

extern __shared__ char shared[]; 
    char data = shared[BaseIndex + tid];