CUDA fundió el acceso a la memoria global

Question

He leído la guía de programación CUDA, pero me faltaba una cosa. Digamos que tengo un array de 32bit int en la memoria global y quiero copiarlo a la memoria compartida con acceso combinado. La matriz global tiene índices de 0 a 1024, y digamos que tengo 4 bloques cada uno con 256 subprocesos.

__shared__ int sData[256]; 

¿Cuándo se realiza el acceso coalescente?

1.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y]; 

Direcciones de memoria global se copian de 0 a 255, cada uno por 32 hilos de urdimbre en, por lo que aquí está bien?

2.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y + someIndex]; 

Si someIndex no es múltiplo de 32 que no se unieron? ¿Direcciones desalineadas? ¿Es eso correcto?

Answer 1

Las normas cuya accesos pueden ser ido formándose son algo complicados y han cambiado con el tiempo. Cada nueva arquitectura CUDA es más flexible en lo que se puede unir. Yo diría que no te preocupes por eso al principio. En cambio, acceda a la memoria de la forma que sea más conveniente y luego vea lo que dice el perfilador CUDA.

Answer 2

Sus ejemplos son correctos si la intención de utilizar una cuadrícula 1D y el hilo de geometría. Creo que la indexación que pretendía usar es [blockIdx.x*blockDim.x + threadIdx.x].

Con # 1, los 32 hilos en una urdimbre ejecutan esa instrucción 'simultáneamente' para que sus solicitudes, que son secuenciales y alineadas a 128B (32 x 4), se combinen en las arquitecturas de Tesla y Fermi, creo.

Con # 2, que es un poco borrosa. Si someIndex es 1, entonces no fusionará todas las 32 solicitudes en un warp, pero podría fusionarse parcialmente. Creo que los dispositivos de Fermi fusionarán los accesos para los subprocesos 1-31 en una distorsión como parte de un segmento secuencial de memoria 128B (y los primeros 4B, que no se necesitan subprocesos, se desperdician). Creo que los dispositivos de arquitectura de Tesla lo convertirían en un acceso no reforzado debido a la desalineación, pero no estoy seguro.

Con someIndex como, por ejemplo, 8, Tesla tendrá 32B direcciones alineadas, y Fermi podría agruparlas como 32B, 64B y 32B. Pero la conclusión es, dependiendo del valor de someIndex y la arquitectura, lo que sucede es borrosa, y no necesariamente será terrible.

Answer 3

Su indexación en 1 es incorrecta (o intencionalmente tan extraña que parece incorrecta), algunos bloques tienen acceso al mismo elemento en cada subproceso, por lo que no hay forma de que estos bloques tengan acceso coalescente.

Prueba:

Ejemplo:

Grid = dim(2,2,0) 

t(blockIdx.x, blockIdx.y) 

//complete block reads at 0 
t(0,0) -> sData[threadIdx.x] = gData[0]; 
//complete block reads at 2 
t(0,1) -> sData[threadIdx.x] = gData[2]; 
//definetly coalesced 
t(1,0) -> sData[threadIdx.x] = gData[threadIdx.x]; 
//not coalesced since 2 is no multiple of a half of the warp size = 16 
t(1,1) -> sData[threadIdx.x] = gData[threadIdx.x + 2]; 

Así que es un juego de "suerte" si un bloque se fundió, lo que en general Sin

Pero coalescencia de memoria lee las reglas son no tan estricto en versiones cuda más nuevas como antes.
Sin embargo, para los problemas de compatibilidad que debe tratar de optimizar los granos para las versiones más bajas cuda, si es posible.

Aquí es alguna fuente agradable:

http://mc.stanford.edu/cgi-bin/images/0/0a/M02_4.pdf

Answer 4

Lo que se quiere en última instancia, depende de si los datos de entrada es una 1D o 2D matriz, y si su red y los bloques son 1D o 2D. El caso más simple es ambos 1D:

shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x]; 

Esto se fusiona. La regla de oro que uso es que la coordenada que varía más rápidamente (el threadIdx) se agrega como compensación al desplazamiento del bloque (blockDim * blockIdx). El resultado final es que la zancada de indexación entre hilos en el bloque es 1. Si la zancada se agranda, entonces pierde la fusión.

La regla simple (en Fermi y GPU posteriores) es que si las direcciones para todos los hilos en un warp caen en el mismo rango alineado de 128 bytes, se producirá una única transacción de memoria (suponiendo que el almacenamiento en caché está habilitado para la carga , que es el valor predeterminado). Si caen en dos rangos alineados de 128 bytes, se producen dos transacciones de memoria, etc.

En GT2xx y GPU anteriores, se vuelve más complicado. Pero puedes encontrar los detalles de eso en la guía de programación.

ejemplos adicionales:

No fusionado:

shmem[threadIdx.x] = gmem[blockDim.x + blockIdx.x * threadIdx.x]; 

No se fundieron, pero no tan malo en GT200 y más tarde:

stride = 2; 
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x]; 

No fusionado en absoluto:

stride = 32; 
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x]; 

Coa lesced, rejilla 2D, bloque 1D:

int elementPitch = blockDim.x * gridDim.x; 
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch + 
          blockIdx.x * blockDim.x + threadIdx.x]; 

coalescentes, rejilla 2D y bloque:

int x = blockIdx.x * blockDim.x + threadIdx.x; 
int y = blockIdx.y * blockDim.y + threadIdx.y; 
int elementPitch = blockDim.x * gridDim.x; 
shmem[threadIdx.y * blockDim.x + threadIdx.x] = gmem[y * elementPitch + x];