Manejo de Memoria en CUDA (I)

1. Manejo de Memoria en CUDA(I) Martín Belzunce Computación Paralela con Procesadores Gráficos Departamento de Electrónica Facultad Regional Buenos Aires – Universidad Tecnológica Nacional Manejo de Memoria en CUDA

2. Temario • Tipos de Memoria • Memoria Global • AccesoUnificado (Coalesced) • CopiasAsincrónicas Host-Device • MemoriaCompartida • Bancos de Memoria • Texturas Manejo de Memoria en CUDA

3. Reglas Generales para Optimizar Código en CUDA • Optimizar Transferencias de Memoria • Maximizar Ocupación del Procesador • Maximizar la Intensidad Aritmética Manejo de Memoria en CUDA

4. Tipos de Memoria en CUDA Memoria Local Privada a cada Thread Memoria Compartida Común a todos los Thread de un mismo Bloque Visible por todos los Threads en ejecución Memoria Global Memoria de Textura Accesible desde el Host Manejo de Memoria en CUDA

5. Tipos de Memoria en las GPUs NVIDIA(Hardware) Memorias Disponibles en Todas las GPUs NVIDIA compatibles con CUDA Manejo de Memoria en CUDA

6. Tipos de Memoria en las GPUs NVIDIA(Hardware) Propiedades de las memorias en las arquitecturas G80 y Tesla T10 (Compute Capabilities 1.x) Manejo de Memoria en CUDA

7. Recordemos Arquitectura Tesla T10 • GPU escalable en cantidad de StreamingMultiprocessors (SMs) • Cada SM contiene: • 8 Procesadores Escalares o ThreadProcessors. • 1 unidad de punto flotante de doble precisión • 2 unidades de ejecución de funciones especiales (SFUs) Manejo de Memoria en CUDA

8. Ámbito de Ejecución de los Threads Manejo de Memoria en CUDA

9. Ejecución de Threads • Los bloques se ejecutan en grupos de 32 Threads denominados Warps. • Los Threads del Warp se distribuyen en los procesadores disponibles dentro del SM. • Acceso a Memoria Global de los Threads en Ejecución de medio Warp • Todos los kernels requieren traer datos de la memoria global • Permite obtener buen ancho de banda (más de 100 Gbytes/sec) • La latencia es muy alta Manejo de Memoria en CUDA

10. Memoria en Arquitectura Fermi • Se introdujeron memorias cache de L1 y L2. • La cache de L1 se comparte con la memoria compartida (configurable). • Ahora la memoria global tiene cache (L2). Manejo de Memoria en CUDA

11. Optimización de Memoria Global • El acceso a los datos se inicia desde la memoria global. (Incluso cuando luego uso memoria compartida) • En la GPU se tienen cientos de cores-> se debe conseguir gran ancho de banda para brindarle constantemente datos. • La mayoría de las aplicaciones están limitadas en ancho de banda -> es fundamental optimizar los accesos a memoria global. (puede enmascara otras optimizaciones) Manejo de Memoria en CUDA

12. Optimización de Memoria Global (II) • El acceso a memoria global tiene una latencia de 400-800 ciclos de clock. • Los threads de un warp quedan frenados cuando requieren un dato de memoria global. • El scheduler ejecuta otro warp durante esta espera. • - > Debo tener muchos warps en ejecución por SM (Alta Ocupación). • Lo deseable es trabajar con el PCI Express saturado (No siempre posible). Manejo de Memoria en CUDA

13. Optimización de Memoria Global (III) • Buena ocupación: +512 threads por SM. • Configurar tamaño de bloque para lograrlo. • Solo hasta 8 bloques por SM. • Bloques chicos - > baja ocupación. • Bloques grandes -> poco flexible. • El tamaño debe ser múltiplo de 32. • Valor de referencia: 128-256 threads. (Hay que evaluar cada aplicación) Manejo de Memoria en CUDA

14. Accesos a Memoria Global (HW) • El ancho del bus depende de la GPU (varios bytes) • Se pueden traer múltiples datos en una única lectura. • Ancho de banda depende además del clock de la memoria. • Ejemplo: clock ddr bus a Bytes BW = 1848.106 [1/seg] * 2 * 384 [bits] / 8 [Bytes/bits] = 177,4 GB/seg Manejo de Memoria en CUDA

15. Accesos a Memoria Global (HW/SW) • ¿Cómo aprovecho todo el ancho de bus? -> Acceso Unificado (coalesced) • Los Threads de un mismo warp acceden a memoria global en la misma lectura (bajo ciertas condiciones) En Tesla de a medio Warp Los Threads deben acceder a posiciones contiguas de memoria El warp podía estar incompleto Manejo de Memoria en CUDA

16. Accesos a Memoria Global (HW/SW) Los accesos a memoria de medio warp se unifican en una o más transacciones dependiendo la alineación y el tamaño de los datos. Para compute capabilities 1.0 y 1.1, sólo coalesced si: Si palabra es de 4 bytes, deben acceder al mismo segmento de 64 bytes. Si es de 8 bytes, deben acceder al mismo segmento de 128 bytes. Threads alineados a memoria. Manejo de Memoria en CUDA

17. Accesos a Memoria Global (HW/SW) • En compute capabilities 1.2 y 1.3 pueden estar en cualquier orden dentro del segmento. • El tamaño del segmento cambia según el tamaño de la palabra (1 byte -> segmento 32 bytes) (2 bytes -> segmento 64 bytes) (4, 8 y 16 bytes -> 128 bytes) • En Fermi (compute capabilities 2.x) se trabaja a nivel de warp completo para datos menores o iguales a 4 bytes (segmento de 128 bytes). • En Fermi se incorporan caches de L1 y L2. • En L1 la línea de cache es de 128 bytes, en L2 de 32 bytes. • En Kepler solo la L2 cachea memoria global. L1 para memoria local. Manejo de Memoria en CUDA

18. Probando el Acceso Coalesced • Implementar un kernel que lea un dato de memoria global, lo incremente y lo escriba en la misma posición. • El tamaño del vector que sea de al menos 10 millones. • Promediar los tiempos de ejecución para 10K corridas. • Evaluar distintos tamaños de datos (unsignedchar, int, float, double) • Realizar acceso alineado y no alineado. • Repetir lo mismo para una matriz almacenada en un array, accediéndola en orden fila-columna y viceversa. Manejo de Memoria en CUDA

19. Como Evitar Accesos NoCoalesced • Si el patrón de acceso a memoria es muy irregular, intentar usar texturas (si es posible). • Si todos los threads acceden a la misma posición de memoria, usar memoria constante. • Para estructura de datos usar estructuras con arrays en vez de arrays de estructuras. • O forzar la alineación de la estructura a 4, 8 o 16 bytes. • O usar el tipo de dato float4. Manejo de Memoria en CUDA

20. Ejemplo de Estructura de Arrays Array de Estructuras Estructuras de Arrays struct Point{ float x; float y; float z;} struct Point{ float x[NUM_THREADS]; float y[NUM_THREADS]; float z[NUM_THREADS];} float3 puntos[NUM_THREADS]; = struct Point puntos[NUM_THREADS]; struct Point puntos; 36 0 12 24 Manejo de Memoria en CUDA 0 32

21. Utilización de Float3 Kernel que usa float3 __global__ voidsimulaParticulas (float3* pos) { intindex = blockIdx.x *blockDim.x+threadIdx.x; float3 localPos = pos[index]; //procesamiento } • Memoria alineada en 12 bytes (distinto a 4, 8 o 16) • Debo hacer 3 lecturas. 1/3 BW Solución 1: Separar en tres variables distintas __global__ voidsimulaParticulas (float* posx, float* posy, float*posz) { intindex = blockIdx.x *blockDim.x+threadIdx.x; float3 localPosx = posx[index]; float3 localPosy= posy[index]; float3 localPosxz= posz[index]; Solución 2: Usar float3 pero cargar a través de floats Manejo de Memoria en CUDA

22. Más Sobre Caches L1 y L2 • La cache L1 comparte la memoria física con la Memoria Compartida. Configurable 16KB L1 / 48KB smem o 48KB L1 / 16 KB L1. • cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferShared); • cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferL1); • L1 sirve para resguardar registros o acceso a datos desalineados. • Las líneas de cache de L1 son de 128 bytes mientras que las de L2 son de 32 bytes. • Que haya un hit en cache de L1 no implica una mejora en el rendimiento (incluso puede empeorar). • Memoria Global con Cache L1 y L2:-Xptxas-dlcm=ca • Solo L2: -Xptxas-dlcm=cg Manejo de Memoria en CUDA

23. Ejemplo Accesos Con y Sin Cache Cada thread del warp pide una palabra de 4 Bytes. 128 Bytes en total. Todas las direcciones caen en una línea de cache. Caching Load: una única transacción. Bus al 100% Non-Caching Load: cache miss. Cuatro transacciones de 32 Bytes. Bus al 100% Manejo de Memoria en CUDA

24. Ejemplo Accesos Con y Sin Cache Las direcciones caen en dos líneas de cache (desalineado). Caching Load: dos transacciones. Bus al 50%. (Se necesitan 128 Bytes, se mueven 256). Non-Caching Load: por estar desalineados caen en 5 segmentos de 32 Bytes. Bus al 80% Las direcciones caen en N en líneas de cache distintas. Caching Load: dos transacciones. Bus al 128/(N*128) % Non-Caching Load: cache miss. Bus al 128/(N*32)% Manejo de Memoria en CUDA

25. Ejemplo Accesos Con y Sin Cache Cada thread del warp pide una palabra de 4 Bytes. 128 Bytes en total. Todas las direcciones caen en una línea de cache. Caching Load: una única transacción. Bus al 100% Non-Caching Load: cache miss. Cuatro transacciones de 32 Bytes. Bus al 100% Manejo de Memoria en CUDA

26. Procesamiento de los Datos • Una vez que fueron cargados los datos de memoria global, se procesan. • Hay que evitar seguir leyendo los datos de memoria global (mucha latencia). • En algunas aplicaciones threads dentro de un mismo bloque deben acceder a los mismos datos. • En estos casos voy a cargar los datos en memoria compartida y manipularlos ahí. Manejo de Memoria en CUDA

27. Utilización de Memoria Compartida Ciclo de Trabajo Datos en Memoria de Host En Memoria de Device (Memoria Global) Se cargan en Memoria Compartida Procesamiento de los Datos Resultados en Memoria Compartida Resultado en Memoria Global Resultados en Memoria de Host Manejo de Memoria en CUDA

28. Implementación de Memoria Compartida (HW) En Tesla (G80 y GT200 o T10) • Memoria Compartida ubicada en el SM. • 16 kBytes de Memoria Compartida por SM. • El espacio se comparte entre todos los bloques ejecutándose en el SM. (HW) • Los datos solo se pueden compartir entre threads del mismo bloque (SW). Manejo de Memoria en CUDA

29. Bancos de Memoria Compartida • Implementada con 16 bancos organizados para que palabras sucesivas de 32 bits accedan a bancos distintos y sucesivos. • Cada banco tiene un ancho de banda de 32 bits / 2 clocks. • Un pedido de acceso a memoria compartida realizado por un warp se divide en dos transacciones. • Si más de un thread quiere acceder al mismo banco se serializan los accesos. (Conflicto de Bancos) • Si no hay conflicto de bancos, rendimiento comparable a registros. • Tiene Broadcast. Manejo de Memoria en CUDA

30. Bancos de Memoria Compartida • No hay conflictos de bancos si cada thread de cada medio warp accede a un banco distinto. Manejo de Memoria en CUDA

31. Conflicto de Bancos • Stride: salto entre datos contiguos. • Si hace que dos threads accedan al mismo dato, se serializan los accesos. Manejo de Memoria en CUDA

32. Conflicto de Bancos (II) • Si tengo los threadstid y tid+n voy a tener conflicto si stride*n es múltiples del nº de bancos. • Si el stride no es divisor de 16, no hay conflictos. • Si voy a leer o escribir datos de 1 o 2 bytes, será muy probable que tenga conflicto de bancos: • __shared__ char datos[NUM_THREADS]; • charaux = datos[Base + threadIdx.x]; • charaux = datos[Base + threadIdx.x*4]; • Para datos mayores a 32 bits voy a tener indefectiblemente conflictos (double o estructuras). • Si todos los threads acceden al mismo banco -> Broadcast-> No hay conflicto. Conflicto Sin Conflicto Manejo de Memoria en CUDA

33. Bancos de Memoria Compartida en Fermi • Memoria compartida con 32 bancos de manera que palabras contiguas de 4 bytes queden en bancos contiguos. • Cada banco 32 bits / 2 clocks. • Ahora si puede haber conflicto entre el primer y el segundo medio warp. • Hay accesos de 64 y 128 bits. No hay más conflictos en array de doubles. • Si tengo estructuras de 128 bits seguramente haya conflicto entre pares de bancos. Manejo de Memoria en CUDA

34. Ejemplos de Acceso a Bancos • Utilización de float3 en Memoria Compartida sin conflicto de bancos: • Acceso a un array 2D __shared__ mat[32][32]; Conflicto de Bancos x16 (Tesla) x32(Fermi) Convierto a __shared__ mat[32][33]; La última columna la lleno con ceros Manejo de Memoria en CUDA

35. Pendientes…. • Memoria de Textura • Memoria Constante • Memoria de Host Pinned • Copias Asincrónicas de Memoria Host-Device Manejo de Memoria en CUDA