Conflicto GPU Banco de Memoria Compartida

Question

Estoy tratando de entender cómo los conflictos bancarios tienen lugar.
si tengo una matriz de tamaño 256 en la memoria global y tengo 256 hilos en un solo bloque, y quiero copiar la matriz de memoria compartida. por lo tanto cada copias de rosca un elemento.

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

hace este resultado simple acción en un conflicto banco?

Supongamos ahora que el tamaño de la matriz es mayor que el número de hilos, por lo que ahora estoy usando 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];

lo hace el resultado código anterior en un conflicto banco?

Answer 1

La mejor manera de comprobar esto sería el perfil de su código utilizando el "Calcular Visual Profiler"; esto viene con el kit de herramientas CUDA. También hay una gran sección en la GPU Gems 3 en esto - "39.2.3 Evitar Banco conflictos".

" 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 4bytes amplia. Así que, esencialmente, si tiene cualquier tema en un medio urdimbre la lectura de la memoria de las mismas 4bytes en un banco de memoria compartida, vas a tener conflictos bancarias y serialización etc.

Aceptar para que el primer ejemplo

En primer lugar permite asumen sus matrices se dicen por ejemplo del tipo int ( una palabra de 32 bits ). Su código guarda estos enteros en la memoria compartida, a través de cualquier medio urdimbre el hilo Kth es salvar al banco de memoria Kth. Así, por ejemplo hilo de 0 de la primera media urdimbre salvará a shared_a[0] que está en el primer banco de memoria, el hilo 1 se ahorrará a shared_a[1], cada medio de urdimbre tiene 16 hilos de estos mapas para los 16 bancos 4byte. En la siguiente media urdimbre, el primer hilo será ahora Guardar su valor en shared_a [16] que está en el banco primero memoria de nuevo. Así que si se utiliza una palabra 4byte tales int, float, etc, entonces el primer ejemplo no dará lugar a un conflicto banco. Si utiliza una palabra 1 byte como el carbón, en los primeros hilos de urdimbre media 0, 1, 2 y 3 será todo Guardar sus valores en el primer banco de la memoria compartida, que causará un conflicto banco.

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 Rosca 1: escribirá a 1, 32

Todos los hilos a través del medio de urdimbre se ejecutan concurrentemente por lo que las escrituras en la memoria compartida no debe causar conflictos bancarios. Voy a tener que vuelva a comprobar éste sin embargo.

Espero que esta ayuda, lo siento por la enorme respuesta!

Answer 2

En ambos casos, el acceso hilos de la memoria compartida con la dirección consecutiva. Depende del tamaño del elemento de memoria compartida, pero el acceso a la memoria compartida consecutiva por una urdimbre de hilos no da lugar a un conflicto banco por "pequeños" tamaños de elementos.

este código con NVIDIA Visual Profiler muestra que para el tamaño del elemento más pequeño que 32 y un múltiplo de 4 (4, 8, 12, ..., 28), el acceso consecutivo a la memoria compartida no da lugar a un conflicto banco. tamaño del elemento de 32, sin embargo, los resultados en conflicto banco.

---

Respuesta por Ljdawson contiene alguna información obsoleta:

... Si utiliza una palabra 1 byte como el carbón, en los primeros hilos de urdimbre media 0, 1, 2 y 3 será todo Guardar sus valores en el primer banco de la memoria compartida, que causará un conflicto banco.

Esto puede ser cierto para las GPU de edad, pero para los últimos GPUs con CC> = 2.x, que no causan conflictos bancarias, de manera efectiva, debido al mecanismo de difusión ( enlace ). Tras la cita es de CUDA C Guía de programación G3.3 (v8.0.61). Memoria compartida .

Una solicitud de memoria para una urdimbre compartida no genera un conflicto banco entre dos hilos que tienen acceso 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 de solicitan (varias palabras se pueden difundir en una sola transacción) y para los accesos de escritura, cada dirección está escrito por uno solo de los hilos (que lleva a cabo del hilo la escritura es indefinido).

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];