2D memoria compartida dispuesta en CUDA

Question

Siempre he trabajado con la memoria lineal compartida (carga, tienda, vecinos de acceso), pero he hecho una prueba sencilla en 2D para estudiar los conflictos bancarias que los resultados me han confundido.

El siguiente código de leer los datos de una matriz de memoria global de dimensiones a la memoria compartida y copiarlo desde la memoria compartida de la memoria global.

__global__ void update(int* gIn, int* gOut, int w) { 

// shared memory space 
__shared__ int shData[16][16]; 
// map from threadIdx/BlockIdx to data position 
int x = threadIdx.x + blockIdx.x * blockDim.x; 
int y = threadIdx.y + blockIdx.y * blockDim.y; 
// calculate the global id into the one dimensional array 
int gid = x + y * w; 

// load shared memory 
shData[threadIdx.x][threadIdx.y] = gIn[gid]; 
// synchronize threads not really needed but keep it for convenience 
__syncthreads(); 
// write data back to global memory 
gOut[gid] = shData[threadIdx.x][threadIdx.y]; 
} 

El perfilador visual informó conflictos de memoria compartida. Los próximos conflictos código de evitar thouse (sólo muestran las diferencias)

// load shared memory 
shData[threadIdx.y][threadIdx.x] = gIn[gid]; 

// write data back to global memory 
gOut[gid] = shData[threadIdx.y][threadIdx.x]; 

Este comportamiento me ha confundido porque en la programación de los procesadores masivamente paralelos. Un enfoque práctico que podemos leer:

Los elementos de matriz en C y CUDA se colocan en las ubicaciones direccionadas linealmente de acuerdo con la convención principal de filas. Es decir, los elementos de la fila 0 de una matriz se colocan primero en orden en ubicaciones consecutivas.

¿Está relacionado con el arragmento de memoria compartida? o con índices de hilos? Tal vez me estoy perdiendo algo?

La configuración del núcleo es el siguiente:

// kernel configuration 
dim3 dimBlock = dim3 (16, 16, 1); 
dim3 dimGrid = dim3 (64, 64); 
// Launching a grid of 64x64 blocks with 16x16 threads -> 1048576 threads 
update<<<dimGrid, dimBlock>>>(d_input, d_output, 1024); 

Gracias de antemano.

Answer 1

Sí, la memoria compartida está ordenada en orden mayor a la esperada. Por lo que su [16] [16] matriz se almacena fila sabia, algo como esto:

 bank0 .... bank15 
row 0 [ 0 .... 15 ] 
    1 [ 16 .... 31 ] 
    2 [ 32 .... 47 ] 
    3 [ 48 .... 63 ] 
    4 [ 64 .... 79 ] 
    5 [ 80 .... 95 ] 
    6 [ 96 .... 111 ] 
    7 [ 112 .... 127 ] 
    8 [ 128 .... 143 ] 
    9 [ 144 .... 159 ] 
    10 [ 160 .... 175 ] 
    11 [ 176 .... 191 ] 
    12 [ 192 .... 207 ] 
    13 [ 208 .... 223 ] 
    14 [ 224 .... 239 ] 
    15 [ 240 .... 255 ] 
     col 0 .... col 15 

Debido a que hay 16 32 bits compartida bancos de memoria sobre pre-Fermi de hardware, cada entrada de número entero en cada columna se asigna a uno compartido Banco de memoria. Entonces, ¿cómo interactúa eso con su elección de esquema de indexación?

Lo que hay que tener en cuenta es que los hilos dentro de un bloque están numerados en el orden principal de la columna (técnicamente la dimensión x de la estructura es la más rápida, seguida de y seguida de z). Así que cuando se utiliza este sistema de indexación:

shData[threadIdx.x][threadIdx.y] 

hilos dentro de una media urdimbre será la lectura de la misma columna, lo que implica la lectura del mismo banco de memoria compartida, y se producirán conflictos bancarios. Cuando se utiliza el esquema contrario:

shData[threadIdx.y][threadIdx.x] 

hilos dentro del mismo medio urdimbre van a leer de la misma fila, lo que implica la lectura de cada uno de los 16 bancos de memoria compartida diferentes, no se producen conflictos.