Monografias.com > Computación > Hardware
Descargar Imprimir Comentar Ver trabajos relacionados

Computación de altas prestaciones con GPUs




Enviado por Pablo Turmero



    Monografias.com

    Índice
    Introducción
    Arquitectura CUDA
    Modelo de ejecución
    Programación
    Rendimiento
    1

    Monografias.com

    Computación de altas prestaciones
    Los grandes supercomputadores son la única alternativa para cierto tipo de aplicaciones
    Sin embargo, hay aplicaciones con menos exigencias computacionales

    GPGPU
    Utilizar GPUs para aplicaciones de propósito general
    Muy atractivo en términos de rendimiento, consumo y coste
    Exclusividad

    2

    Monografias.com

    Introducción
    Rendimiento
    Fuente: Nvidia
    3

    Monografias.com

    Introducción
    Precios
    Más de 100 millones de GPUs de Nvidia
    4

    Monografias.com

    Introducción
    Uso de los transistores
    + Cache
    + Control
    – Cores
    + Cores
    – Cache
    – Control
    5

    Monografias.com

    Introducción
    6

    Monografias.com

    Introducción
    GeForce 8800 GTX vs. 2.2GHz Opteron 248
    10? speedup en un kernel es típico, si hay suficiente paralelismo
    25? a 400? speedup con optimizaciones
    Depende mucho del tipo de aplicación
    7

    Monografias.com

    Introducción
    Programabilidad
    2005
    2009
    Ensamblador
    OpenGL
    DirectX
    Cg
    Brook++
    RapidMind
    CUDA
    OpenCL
    Facilidad de programación
    Inicios GPGPU
    ATI Stream
    Explosión GPGPU
    Futuro GPGPU?
    8

    Monografias.com

    Índice
    Introducción
    Arquitectura CUDA
    Modelo de ejecución
    Programación
    Rendimiento
    9

    Monografias.com

    Arquitectura CUDA
    . . .
    Memoria global
    Conjunto de Streaming Multiprocessors (MP)
    (Gp:) SP
    (Gp:) DP
    (Gp:) SM

    8 Scalar Processors (SP)
    1 Unidad de Doble Precisión (DP)
    16 KB de Memoria Compartida (SM)
    8-16 K Registros
    10

    Monografias.com

    Fermi
    ~1.5TFLOPS (SP)
    ~800GFLOPS (DP)
    230 GB/s DRAM
    11

    Monografias.com

    Espacios de memoria en CUDA
    12
    Fuente: Nvidia

    Monografias.com

    Espacios de memoria en CUDA
    13

    Monografias.com

    Gestión de memoria con CUDA
    CPU y GPU tienen espacios de memoria independientes
    Transferencia de datos entre ambos espacios de memoria a través del bus PCIExpress
    Reserva, transferencia y liberación explícitas
    Las operaciones con memoria realizadas por el host
    14
    Fuente: Nvidia

    Monografias.com

    Gestión de memoria con CUDA
    cudaMalloc()
    Obtiene espacio en la memoria global
    Parámetros: dirección del puntero y el tamaño a reservar
    cudaMemset()
    Inicializa a un valor dado
    Parámetros: dirección, valor y cantidad
    cudaFree()
    Libera el espacio
    Parámetros: dirección
    15
    Fuente: Nvidia

    Monografias.com

    Gestión de memoria con CUDA
    cudaMemcpy()
    Transfiere datos entre memorias
    Requiere 4 parámetros:
    Puntero al destino
    Puntero al origen
    Bytes a copiar
    Tipo de transferencia:
    Host a host
    Host a dispositivo
    Dispositivo a host
    Dispositivo a dispositivo
    16
    Fuente: Nvidia

    Monografias.com

    Ejemplo
    Reservar, inicializar y liberar espacio para una matriz 64×64 de elementos float
    Enlazar ese espacio a Md
    Gestión de memoria con CUDA
    #define BLOCK_SIZE 64
    float * Md;
    int size = BLOCK_SIZE*BLOCK_SIZE*sizeof(float);
    cudaMalloc((void**)&Md,size);
    cudaMemset(Md,0,size);
    . . .
    cudaFree(Md);
    17

    Monografias.com

    Ejemplo
    Transferir una matriz de 64×64 float de la memoria de la GPU a la del host y viceversa
    Matriz M está en el host y la matriz Md en el dispositivo
    Gestión de memoria con CUDA
    cudaMemcpy(M,Md,size,cudaMemcpyDeviceToHost);

    cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice);
    18

    Monografias.com

    Índice
    Introducción
    Arquitectura CUDA
    Modelo de ejecución
    Programación
    Rendimiento
    19

    Monografias.com

    Modelo de ejecución
    Secuencial
    Código
    Thread
    20

    Monografias.com

    Modelo de ejecución
    Paralelo
    Código
    Threads
    21

    Monografias.com

    kernel
    Paralelo
    Miles de threads
    lanzados a la vez
    22

    Monografias.com

    Modelo de ejecución
    Memoria
    Kernel
    Threads
    (instancias del
    kernel)
    PCIe
    CPU
    GPU
    23

    Monografias.com

    Modelo de ejecución
    CPU
    GPU
    Cada thread tiene un identificador
    Todos los threads ejecutan el mismo código
    Cada thread opera sobre distintos datos

    Modelo SIMD
    CPU

    Threads pesados
    Sobrecarga planificación
    Cambios de contexto lentos
    CPU

    Threads ligeros
    Poca sobrecarga planificación
    Cambios de contexto rápidos
    24

    Monografias.com

    Estructura jerárquica de threads
    Los threads se agrupan en bloques de threads
    Los bloques de threads se agrupan en Grids
    Los threads de un bloque se comunican mediante la memoria compartida
    Todos los threads del Grid se comunican a través de la memoria global

    Un thread se ejecuta en un procesador escalar (SP)
    Un bloque de threads se lanza en un multiprocesador (MP)
    Un MP puede desarrollar varios bloques
    (Gp:) Grid 1
    (Gp:) Bloque
    (2,1)
    (Gp:) Bloque
    (2,0)
    (Gp:) Bloque
    (0,0)
    (Gp:) Bloque
    (1,0)
    (Gp:) Bloque
    (0,1)
    (Gp:) Bloque
    (1,1)

    (Gp:) Thread
    (0,0)
    (Gp:) Thread
    (1,0)
    (Gp:) Thread
    (2,0)
    (Gp:) Thread
    (3,0)
    (Gp:) Thread
    (4,0)
    (Gp:) Thread
    (0,1)
    (Gp:) Thread
    (1,1)
    (Gp:) Thread
    (2,1)
    (Gp:) Thread
    (3,1)
    (Gp:) Thread
    (4,1)
    (Gp:) Thread
    (0,2)
    (Gp:) Thread
    (1,2)
    (Gp:) Thread
    (2,2)
    (Gp:) Thread
    (3,2)
    (Gp:) Thread
    (4,2)
    (Gp:) Bloque (1,1)

    25

    Monografias.com

    Modelo de ejecución
    Software
    Hardware
    (Gp:) Thread
    (Gp:) SP

    (Gp:) Un bloque se ejecuta en un MP
    Los bloques no migran entre MPs
    Varios bloques a la vez en un MP
    (según registros y memoria compartida)

    (Gp:) Bloque threads
    (Gp:) MP

    (Gp:) …
    (Gp:) Grid
    (Gp:) GPU
    (Gp:) Un único kernel concurrente

    Cada thread se ejecuta en un SP
    26

    Monografias.com

    Indentificadores y dimensiones
    El tamaño del grid y de los bloques los determina el programador
    Se usan las variables gridDim y blockDim para referenciar la dimensión de grid y bloque, respectivamente
    (Gp:) Grid 1
    (Gp:) Bloque
    (2,1)
    (Gp:) Bloque
    (2,0)
    (Gp:) Bloque
    (0,0)
    (Gp:) Bloque
    (1,0)
    (Gp:) Bloque
    (0,1)
    (Gp:) Bloque
    (1,1)

    (Gp:) Thread
    (0,0)
    (Gp:) Thread
    (1,0)
    (Gp:) Thread
    (2,0)
    (Gp:) Thread
    (3,0)
    (Gp:) Thread
    (4,0)
    (Gp:) Thread
    (0,1)
    (Gp:) Thread
    (1,1)
    (Gp:) Thread
    (2,1)
    (Gp:) Thread
    (3,1)
    (Gp:) Thread
    (4,1)
    (Gp:) Thread
    (0,2)
    (Gp:) Thread
    (1,2)
    (Gp:) Thread
    (2,2)
    (Gp:) Thread
    (3,2)
    (Gp:) Thread
    (4,2)
    (Gp:) Bloque (1,1)

    (Gp:) gridDim.x
    (Gp:) gridDim.y

    Un thread queda indentificado por:
    Un identificador propio dentro del bloque al que pertenece
    El identificador del bloque al que pertenece
    Se usan las variables threadIdx y blockIdx para referenciar el identificador del thread dentro del bloque y al bloque dentro del grid, respectivamente
    (Gp:) blockDim.x
    (Gp:) blockDim.y
    (Gp:) blockDim.z

    (Gp:) blockIdx.x=1
    blockIdx.y=1

    (Gp:) threadIdx.x=0
    threadIdx.y=2

    27

    Monografias.com

    Planificación
    Se agrupan los threads en bloques
    Se asignan identificadores a bloques y threads
    Se distribuyen los bloques de threads entre los multiprocesadores
    Los threads de un bloque se ejecutan concurrentemente en un multiprocesador
    (Gp:) tiempo
    (Gp:) warp 8 instrucción 11
    (Gp:) warp 1 instrucción 42
    (Gp:) warp 3 instrucción 95
    (Gp:) warp 8 instrucción 12
    (Gp:) …
    (Gp:) warp 3 instrucción 96

    (Gp:) Bloque 1
    (Gp:) Bloque 2
    (Gp:) Bloque n
    (Gp:) warp 1
    (Gp:) 2
    (Gp:) m
    (Gp:) warp 1
    (Gp:) 2
    (Gp:) m
    (Gp:) warp 1
    (Gp:) 2
    (Gp:) m

    Los threads de un bloque son agrupados en warps
    Un warp es la unidad mínima de planificación y está formada por 32 threads
    Varios warps en cada multiprocesador, pero sólo uno está en ejecución en cada momento
    Los warps cuyos operandos están listos son seleccionados para ejecución
    Todos los threads en un warp ejecutan la misma instrucción
    28

    Monografias.com

    Planificación
    Tres flujos de instrucciones: warp1, warp3 y warp8
    t=k
    t=k+1
    t=k+2
    t=l>k
    t=l+1
    Planifica
    en tiempo k
    (Gp:) warp 8 instrucción 11

    (Gp:) warp 1 instrucción 42

    (Gp:) warp 3 instrucción 95

    (Gp:) warp 8 instrucción 12


    (Gp:) warp 3 instrucción 96

    29

    Monografias.com

    Planificación
    (Gp:) warp 8 instrucción 11

    (Gp:) warp 1 instrucción 42

    (Gp:) warp 3 instrucción 95

    (Gp:) warp 8 instrucción 12


    (Gp:) warp 3 instrucción 96

    t=k
    t=k+1
    t=k+2
    t=l>k
    t=l+1
    Planifica en
    tiempo k+1
    Tres flujos de instrucciones: warp1, warp3 y warp8
    30

    Monografias.com

    Índice
    Introducción
    Arquitectura CUDA
    Modelo de ejecución
    Programación
    Rendimiento
    31

    Monografias.com

    Programación
    Computación heterogénea
    Parte de código se ejecuta en la CPU
    Parte de código se ejecuta en la GPU (kernel)

    La API de CUDA es una extensión al lenguaje ANSI C
    Curva de aprendizaje suave

    Extensiones básicas
    Modificadores de función
    Modificadores de variables
    Variables específicas de dimensionado
    Directiva para ejecución del kernel
    32

    Monografias.com

    Modificadores de función
    Indica dónde se ejecuta la función: GPU (device) o CPU (host)

    __device__ La función debe ejecutarse en el dispositivo
    Sólo puede ser llamada por el propio dispositivo
    Recursividad no soportada
    No pueden declararse variables estáticas dentro de la función
    La función no puede tener un número variable de argumentos
    __global__ La función es un kernel que debe ejecutarse en el dispositivo
    Sólo puede ser llamada por el host
    Recursividad no soportada
    No pueden declararse variables estáticas dentro de la función
    La función no puede tener un número variable de argumentos
    La función debe devolver siempre void
    __host__ La función debe ejecutarse en el host
    Sólo puede ser llamada por el host
    No puede utilizarse junto con __global__
    33

    Monografias.com

    Modificadores de variables
    Indica en qué parte de la memoria se localiza la variable

    __device__ La variable reside en el dispositivo
    Requiere que se indique uno de los otros dos modificadores de variables para indicar dónde exactamente reside la variable en el dispositivo
    __constant__ La variable reside en el espacio de memoria constante del dispositivo
    Está viva durante todo el tiempo de ejecución de la aplicación
    Accesible por todos los threads del grid, así como desde el host
    __shared__ La variable reside en el espacio de memoria compartida del bloque de threads en el dispositivo
    Está viva mientras el bloque está vivo
    Accesible sólo por los threads del bloque
    34

    Monografias.com

    Variables específicas dimensiones
    Indican las dimensiones que caracterizan a los identificadores de los threads y bloques

    dim3 gridDim;
    Dimensiones de los grids en bloques (gridDim.z no usado)
    dim3 blockDim;
    Dimensiones del bloque en threads
    dim3 blockIdx;
    Indice del bloque en el grid
    dim3 threadIdx;
    Indice del thread en el bloque

    35

    Monografias.com

    Directiva ejecución del kernel
    Indica cómo debe ejecutarse el kernel en el dispositivo

    Cada vez que se llama a __global__, también debe especificarse la configuración de ejecución del kernel
    Se inserta entre el nombre de la función y el paréntesis del argumento una expresión de la forma: >
    Dg: indica las dimensiones de la malla, es decir, la cantidad de bloques
    Db: indica las dimensiones de cada bloque, es decir, el número de threads por bloque
    Ns: indica el número de bytes de memoria compartida asignada dinámicamente por bloque (argumento opcional, 0 por defecto)
    S: especifica un stream (argumento opcional, 0 por defecto)
    Las invocaciones de kernels son asíncronas
    El control vuelve al programa principal (CPU)

    36

    Monografias.com

    Ejecución del kernel
    Un kernel debe ser ejecutado del siguiente modo:

    __global__ void KernelFunction(…);

    dim3 DimGrid(100,50); // 5000 bloques
    dim3 DimBlock(4,8,8); // 256 threads/bloque

    KernelFunction>(…);
    37

    Monografias.com

    Ejemplo: SAXPY
    // Definición de la función
    void saxpy_serial (int n, float a, float *x, float *y)
    {
    for (int i = 0; i < n; ++i)
    y[i] = a*x[i] + y[i];
    }

    // Llamada a la función
    saxpy(n,2.0,x,y)
    // Definición del kernel
    __global__ void saxpy_parallel (int n, float a, float *x, float *y)
    {
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < n) y[i] = a*x[i] + y[i];
    }

    // Llamada al kernel
    int nblocks = (n + 255) / 256;
    saxpy_parallel(n, 2.0, x, y);
    38

    Monografias.com

    Índice
    Introducción
    Arquitectura CUDA
    Modelo de ejecución
    Programación
    Rendimiento
    39

    Monografias.com

    Optimizaciones
    Reducir sobrecargas en el acceso a memoria
    Realizar transferencias asíncronas
    Mejorar accesos a memoria global
    Usar memoria compartida
    Reducir conflictos de bancos en memoria compartida

    Configurar la ejecución
    Seleccionar el número de bloques y threads por bloque

    Optimizar a nivel de instrucción
    Instrucciones aritméticas
    Instrucciones para acceso a memoria
    Saltos y sentencias condicionales

    40

    Monografias.com

    Transferencias asíncronas
    cudaMemcpy es bloqueante
    Se devuelve el control al host una vez finalizada la transferencia

    cudaMemcpyAsync es no bloqueante
    Se devuelve el control inmediatamente

    Las transferencias no bloqueantes permiten solapar computación y comunicación
    41

    Monografias.com

    Solapar computación y comunicación
    Un stream es un conjunto de operaciones que se completan secuencialmente
    Se puede solapar la comunicación y computación de diferentes streams
    Para ello se debe permitir que la memoria del host esté mapeada en el espacio de direcciones del dispositivo
    42
    Datos2
    Datos1
    Datos2
    Datos1
    Host?Device
    Device?Host
    Kernel Cálculos
    Datos1
    Datos2
    Datos2
    Host?Device
    Device ? Host
    Kernel Cálculos
    Datos1
    Datos2
    Datos1
    Datos2
    Datos1
    42

    Monografias.com

    Memoria global
    Latencia alta
    Usarla lo menos posible
    Accesos por half-warp (16 threads)
    Intentar completarlos en el menor número de transacciones posible (coalescing)
    1 transacción
    16 transacciones
    43

    Monografias.com

    Memoria global
    Latencia alta
    Usarla lo menos posible
    Accesos por half-warp (16 threads)
    Intentar completarlos en el menor número de transacciones posible (coalescing)
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    1
    2
    3
    4
    5
    6
    7
    8

    memoria global
    matriz
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    1
    2
    3
    4
    5
    6
    7
    8

    memoria global
    matriz
    44

    Monografias.com

    Memoria compartida
    Dividida en módulos (bancos)
    En las GPUs actuales hay 16 bancos de 1KB
    Acceso simultáneo a los bancos
    Baja latencia (similar registros)
    Problema: conflictos en los bancos (acceso al mismo banco por dos o más threads)
    (Gp:) Banco 0
    (Gp:) Banco 1
    (Gp:) Banco 2
    (Gp:) Banco 3
    (Gp:) Banco 4
    (Gp:) Banco 14
    (Gp:) Banco 15
    (Gp:) Thread 0
    (Gp:) Thread 1
    (Gp:) Thread 2
    (Gp:) Thread 3
    (Gp:) Thread 4
    (Gp:) Thread 14
    (Gp:) Thread 15
    (Gp:) . . .
    (Gp:) . . .
    (Gp:) Sin conflictos
    (Gp:) Direccionamiento lineal
    Stride = 1

    (Gp:) Banco 0
    (Gp:) Banco 1
    (Gp:) Banco 2
    (Gp:) Banco 3
    (Gp:) Banco 4
    (Gp:) Banco 14
    (Gp:) Banco 15
    (Gp:) Thread 0
    (Gp:) Thread 1
    (Gp:) Thread 2
    (Gp:) Thread 3
    (Gp:) Thread 4
    (Gp:) Thread 14
    (Gp:) Thread 15
    (Gp:) . . .
    (Gp:) . . .
    (Gp:) Sin conflictos
    (Gp:) Direccionamiento aleatorio
    Permutación 1:1

    (Gp:) Banco 0
    (Gp:) Banco 1
    (Gp:) Banco 2
    (Gp:) Banco 3
    (Gp:) Banco 4
    (Gp:) Banco 5
    (Gp:) Banco 15
    (Gp:) Thread 0
    (Gp:) Thread 1
    (Gp:) Thread 2
    (Gp:) Thread 3
    (Gp:) Thread 8
    (Gp:) Thread 15
    (Gp:) . . .
    (Gp:) . . .
    (Gp:) . . .
    (Gp:) Thread 9
    (Gp:) Con conflictos
    (Gp:) Direccionamiento lineal
    Stride = 2

    45

    Monografias.com

    Memoria compartida
    Dividida en módulos (bancos)
    En las GPUs actuales hay 16 bancos de 1KB
    Acceso simultáneo a los bancos
    Baja latencia (similar registros)
    Problema: conflictos en los bancos (acceso al mismo banco por dos o más threads)
    Cargar los datos desde la memoria global a la memoria compartida
    Sincronizar threads
    Procesar usando sólo memoria compartida
    Sincronizar
    Llevar resultados a memoria global
    46

    Monografias.com

    Bloques y threads
    Ocupación:
    Número de warps activos por MP con respecto al máximo posible de warps activos

    El número total de bloques y de threads por bloque son factores importantes
    Número de bloques
    Dependiente de los recursos disponibles
    Al menos tantos como MPs
    Más de uno por MP para que unos mantengan los recursos ocupados mientras otros se sincronizan
    Suficientemente alto para escalar en futuras versiones
    Número de threads por bloque
    Múltiplo del tamaño del warp para facilitar coalescing
    Múltiplos de 64, para evitar conflictos en acceso a registros
    Entre 128 y 256 buena elección
    Más threads no implica necesariamente mayor ocupación
    Experimentación
    47

    Monografias.com

    Rendimiento a nivel de instrucción
    La productividad a nivel de instrucción depende de
    Número de operaciones
    Latencia de la memoria
    Ancho de banda de la memoria
    Maximizar ancho de banda efectivo
    Maximizar el uso de la memoria compartida
    Minimizar los accesos a memoria global
    Maximizar coalescing en los accesos a memoria global
    Solapar comunicación y computación
    Aumentar el rendimiento de la ejecución de las instrucciones
    Instrucciones aritméticas más rápidas, si se prefiere velocidad en lugar de precisión
    Instrucciones para accesos a memoria con menos latencia
    Evitar sentencias condicionales que generen diferentes caminos en el mismo warp, pues éstos son serializados
    48

    Monografias.com

    Divergencia de caminos
    (Gp:) Branch
    (Gp:) Path A
    (Gp:) Path B

    Salto
    Path A
    Path B
    Los caminos son serializados
    Incremento del número de instrucciones ejecutadas por el warp
    Los threads vuelven a converger cuando todos los caminos se completan

    50% de pérdida de rendimiento
    49

    Monografias.com

    Índice
    Producto Matriz-Matriz
    ¿Cómo explotar más la arquitectura?
    Buenas prácticas de programación CUDA
    OpenCL
    50

    Monografias.com

    Índice
    Producto Matriz-Matriz
    Implementación secuencial
    Recordar
    Implementación básica
    ¿Cómo compilar?
    Implementación memoria compartida
    ¿Cómo explotar más la arquitectura?
    Buenas prácticas de programación CUDA
    OpenCL
    51

    Monografias.com

    Producto Matriz-Matriz
    Implementación secuencial
    52

    Monografias.com

    Producto Matriz-Matriz
    cudaMalloc()
    cudaMemcpy()
    cudaThreadSynchronize();
    cudaFree()
    multiplication>(d_P, d_M, d_N, M.width, N.width);
    __global__ void multiplication(float *P, float *M, float *N, int wM, int wN)

    53

    Monografias.com

    Producto Matriz-Matriz
    Implementación básica
    Cada thread calcula un elemento de la matriz
    Lee una fila de A
    Lee una columna de B

    54

    Monografias.com

    Producto Matriz-Matriz
    Implementación básica
    Ejercicio 1: (En programa principal)
    a) Repasar las llamadas y el significado
    b) Completar llamada a kernel
    c) Cálculo de GigaFlops
    Ejercicio 2: (En kernel)
    a) Cálculo de índices
    b) Cálculo de elemento
    Ejercicio 3: (Pruebas)

    55

    Monografias.com

    Producto Matriz-Matriz
    Implementación básica
    Ejercicio 2.a (Ejemplo)

    56
    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (2,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (2,1)
    (Gp:) (0,2)
    (Gp:) (1,2)
    (Gp:) (2,2)

    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (2,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (2,1)
    (Gp:) (0,2)
    (Gp:) (1,2)
    (Gp:) (2,2)

    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (2,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (2,1)
    (Gp:) (0,2)
    (Gp:) (1,2)
    (Gp:) (2,2)

    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (2,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (2,1)
    (Gp:) (0,2)
    (Gp:) (1,2)
    (Gp:) (2,2)

    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (2,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (2,1)
    (Gp:) (0,2)
    (Gp:) (1,2)
    (Gp:) (2,2)

    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (2,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (2,1)
    (Gp:) (0,2)
    (Gp:) (1,2)
    (Gp:) (2,2)

    5
    3
    (Gp:) (0,0)
    (Gp:) (1,0)
    (Gp:) (0,1)
    (Gp:) (1,1)
    (Gp:) (0,2)
    (Gp:) (1,2)

    Monografias.com

    Producto Matriz-Matriz
    ¿Cómo compilar?
    gmultiply: gmultiply.cu
    nvcc gmultiply.cu
    -keep
    –ptxas-options=-v
    -o gmultiply
    -lcutil
    -L /usr/local/cuda/sdk/lib
    -I /usr/local/cuda/sdk/common/inc
    57

    Monografias.com

    GPU GeForce 285GTX
    30 multiprocesadores (MPs)
    8 cores (FPUs) por multiprocesador, a 1476 MHz
    mad: multiplicación-suma (considera 2 flops a efectos de cálculo de productividad)
    2 SFU (unidades funcionales especiales) por multiprocesador

    Total 240 cores

    797 GFLOPS (rendimiento pico)
    (30 MPs) x (8×2 + 2 Flops/MP) x 1476 MHz = 797 GFLOPS

    1GB RAM
    1242 MHz (double data rate, 2484 MHz)
    512 bits (ancho bus)
    159 GB/s = (2484 MHz) x (512 / 8 bytes)

    58

    Monografias.com

    Producto Matriz-Matriz
    Se usan los cores (no las SFUs)
    (240 cores) x 1476 MHz =
    = 354,24 Goperaciones/s

    De 8 instrucciones sólo 2 son FLOP
    Máxima productividad
    354,24 x ¼ = 88,56 GLOPS

    ¿Por qué no se llega a ese valor?
    59
    $Lt_0_7:
    ld.global.f32 %f2, [%r15+0];
    ld.global.f32 %f3, [%r17+0];
    mad.f32 %f1, %f2, %f3, %f1;
    add.s32 %r17, %r17, 4096;
    add.s32 %r15, %r15, 4;
    setp.ne.u32 %p1, %r15, %r16;
    @%p1 bra $Lt_0_7;

    Accesos a memoria
    ¼ operaciones son cargas
    (240 cores) x (¼ cargas) x (4bytes/carga) x (1476 MHz) = 354 GBs > 159 GBs

    El ancho de banda con memoria global no es suficiente

    Monografias.com

    Producto Matriz-Matriz
    Cada dato de entrada es leído por width threads

    Latencia memoria global ~ 400 ciclos
    Latencia memoria compartida ~ 10 ciclos

    Llevar cada dato de entrada a memoria compartida y que sea leído por varios threads

    60

    Monografias.com

    Producto Matriz-Matriz
    Cada dato de entrada es leído por width threads

    Latencia memoria global ~ 400 ciclos
    Latencia memoria compartida ~ 10 ciclos

    Llevar cada dato de entrada a memoria compartida y que sea leído por varios threads

    Producto de sub-matrices

    Cada sub-matriz es calculada por un bloque de threads
    Los datos de entrada necesarios son llevados a memoria compartida
    61

    Monografias.com

    Producto Matriz-Matriz
    Implementación memoria compartida
    Disminuir accesos a memoria
    ¿Cómo lo haríais?
    62

    Monografias.com

    Producto Matriz-Matriz
    63
    (Gp:) C1,0
    (Gp:) C0,0
    (Gp:) C0,1
    (Gp:) C2,0
    (Gp:) C3,0
    (Gp:) C1,1
    (Gp:) C0,2
    (Gp:) C2,2
    (Gp:) C3,2
    (Gp:) C1,2
    (Gp:) C3,1
    (Gp:) C2,1
    (Gp:) C0,3
    (Gp:) C2,3
    (Gp:) C3,3
    (Gp:) C1,3

    (Gp:) A2,0
    (Gp:) A1,1
    (Gp:) A1,0
    (Gp:) A0,0
    (Gp:) A0,1
    (Gp:) A3,0
    (Gp:) A2,1
    (Gp:) A3,1
    (Gp:) A2,2
    (Gp:) A1,3
    (Gp:) A1,2
    (Gp:) A0,2
    (Gp:) A0,3
    (Gp:) A3,2
    (Gp:) A2,3
    (Gp:) A3,3

    (Gp:) B0,3
    (Gp:) B1,3
    (Gp:) B1,2
    (Gp:) B1,1
    (Gp:) B1,0
    (Gp:) B0,0
    (Gp:) B0,1
    (Gp:) B0,2
    (Gp:) B2,3
    (Gp:) B3,3
    (Gp:) B3,2
    (Gp:) B3,1
    (Gp:) B3,0
    (Gp:) B2,0
    (Gp:) B2,1
    (Gp:) B2,2

    Monografias.com

    Producto Matriz-Matriz
    64
    orden accesos
    En una primera fase se llevan a memoria compartida
    En una segunda fase se llevan a memoria compartida
    Cada Ai,j y Bi,j se lee dos veces
    Cada thread calcula un punto de la matriz resultado

    Monografias.com

    Producto Matriz-Matriz
    65
    C1,0
    C0,0
    C0,1
    C2,0
    C3,0
    C1,1
    C0,2
    C2,2
    C3,2
    C1,2
    C3,1
    C2,1
    C0,3
    C2,3
    C3,3
    C1,3
    (Gp:) A2,0
    (Gp:) A1,1
    (Gp:) A1,0
    (Gp:) A0,0
    (Gp:) A0,1
    (Gp:) A3,0
    (Gp:) A2,1
    (Gp:) A3,1
    (Gp:) A2,2
    (Gp:) A1,3
    (Gp:) A1,2
    (Gp:) A0,2
    (Gp:) A0,3
    (Gp:) A3,2
    (Gp:) A2,3
    (Gp:) A3,3

    (Gp:) B0,3
    (Gp:) B1,3
    (Gp:) B1,2
    (Gp:) B1,1
    (Gp:) B1,0
    (Gp:) B0,0
    (Gp:) B0,1
    (Gp:) B0,2
    (Gp:) B2,3
    (Gp:) B3,3
    (Gp:) B3,2
    (Gp:) B3,1
    (Gp:) B3,0
    (Gp:) B2,0
    (Gp:) B2,1
    (Gp:) B2,2

    (1,1)
    (1,0)
    (0,0)
    (0,1)
    threads
    C0,0
    C1,0
    C1,1
    C0,1
    calculan
    memoria compartida
    As
    Bs

    cada thread lleva un dato de A y B

    Monografias.com

    Producto Matriz-Matriz
    66
    C1,0
    C0,0
    C0,1
    C2,0
    C3,0
    C1,1
    C0,2
    C2,2
    C3,2
    C1,2
    C3,1
    C2,1
    C0,3
    C2,3
    C3,3
    C1,3
    (Gp:) A2,0
    (Gp:) A1,1
    (Gp:) A1,0
    (Gp:) A0,0
    (Gp:) A0,1
    (Gp:) A3,0
    (Gp:) A2,1
    (Gp:) A3,1
    (Gp:) A2,2
    (Gp:) A1,3
    (Gp:) A1,2
    (Gp:) A0,2
    (Gp:) A0,3
    (Gp:) A3,2
    (Gp:) A2,3
    (Gp:) A3,3

    (Gp:) B0,3
    (Gp:) B1,3
    (Gp:) B1,2
    (Gp:) B1,1
    (Gp:) B1,0
    (Gp:) B0,0
    (Gp:) B0,1
    (Gp:) B0,2
    (Gp:) B2,3
    (Gp:) B3,3
    (Gp:) B3,2
    (Gp:) B3,1
    (Gp:) B3,0
    (Gp:) B2,0
    (Gp:) B2,1
    (Gp:) B2,2

    (1,1)
    (1,0)
    (0,0)
    (0,1)
    threads
    C0,0
    C1,0
    C1,1
    C0,1
    calculan
    memoria compartida
    As
    Bs
    A0,0
    A1,0
    A1,1
    A0,1
    B0,0
    B1,0
    B1,1
    B0,1
    C0,0 = A0,0 x B0,0 + A1,0 x B0,1
    C1,0 = A0,0 x B1,0 + A1,0 x B1,1
    C0,1 = A0,1 x B0,0 + A1,1 x B0,1
    C1,1 = A0,1 x B1,0 + A1,1 x B1,1

    Monografias.com

    Producto Matriz-Matriz
    67
    C1,0
    C0,0
    C0,1
    C2,0
    C3,0
    C1,1
    C0,2
    C2,2
    C3,2
    C1,2
    C3,1
    C2,1
    C0,3
    C2,3
    C3,3
    C1,3
    (Gp:) A2,0
    (Gp:) A1,1
    (Gp:) A1,0
    (Gp:) A0,0
    (Gp:) A0,1
    (Gp:) A3,0
    (Gp:) A2,1
    (Gp:) A3,1
    (Gp:) A2,2
    (Gp:) A1,3
    (Gp:) A1,2
    (Gp:) A0,2
    (Gp:) A0,3
    (Gp:) A3,2
    (Gp:) A2,3
    (Gp:) A3,3

    (Gp:) B0,3
    (Gp:) B1,3
    (Gp:) B1,2
    (Gp:) B1,1
    (Gp:) B1,0
    (Gp:) B0,0
    (Gp:) B0,1
    (Gp:) B0,2
    (Gp:) B2,3
    (Gp:) B3,3
    (Gp:) B3,2
    (Gp:) B3,1
    (Gp:) B3,0
    (Gp:) B2,0
    (Gp:) B2,1
    (Gp:) B2,2

    (1,1)
    (1,0)
    (0,0)
    (0,1)
    threads
    C0,0
    C1,0
    C1,1
    C0,1
    calculan
    memoria compartida
    As
    Bs
    A2,0
    A3,0
    A3,1
    A2,1
    B0,2
    B1,2
    B1,3
    B0,3
    C0,0 = A0,0 x B0,0 + A1,0 x B0,1 + A2,0 x B0,2 + A3,0 x B0,3
    C1,0 = A0,0 x B1,0 + A1,0 x B1,1 + A2,0 x B1,2 + A3,0 x B1,3
    C0,1 = A0,1 x B0,0 + A1,1 x B0,1 + A2,1 x B0,2 + A3,1 x B0,3
    C1,1 = A0,1 x B1,0 + A1,1 x B1,1 + A2,1 x B1,2 + A3,1 x B1,3

    Monografias.com

    Producto Matriz-Matriz
    Bloque de threads ? 16 x 16 = 256 threads
    Width = 4096 ? 256 x 256 = 65536 bloques de threads

    Cada dato en una sub-matriz es leído por 16 threads
    Los accesos a memoria global se reducen en un factor 16 al usar memoria compartida
    Se requiere ahora 354 / 16 ~ 22 GBs

    Ahora la memoria no es motivo para no alcanzar la productividad deseada

    Monografias.com

    Producto Matriz-Matriz
    Implementación memoria compartida
    Ejercicio 1: (En programa principal)
    a) Repasar las llamadas y el significado
    b) Completar llamada a kernel
    c) Cálculo de GigaFlops
    Ejercicio 2: (En kernel)
    a) Cálculo de índices
    b) Cálculo de elemento
    Ejercicio 3: (Pruebas)

    69

    Monografias.com

    Producto Matriz-Matriz
    Implementación memoria compartida
    Disminuir accesos a memoria
    ¿Cómo lo haríais?
    Veamos el código ahora e identifiquemos los diferentes elementos
    70

    Monografias.com

    Índice
    Producto Matriz-Matriz
    ¿Cómo explotar más la arquitectura?
    Buenas prácticas de programación CUDA
    OpenCL
    71

    Monografias.com

    ¿Cómo explotar más la arquitectura?
    Sistemas Distribuidos mediante MPI

    Sistemas compartidos mediante OpenMP
    72

    Monografias.com

    Índice
    Producto Matriz-Matriz
    ¿Cómo explotar más la arquitectura?
    Buenas prácticas de programación CUDA
    OpenCL
    73

    Monografias.com

    Guía de buenas prácticas de programación CUDA
    Maximizar ejecución paralela
    Optimizar el uso de la memoria de cara obtener un mayor ancho de banda en el acceso a memoria
    Optimizar el uso de las instrucciones para conseguir una mayor productividad
    74

    Monografias.com

    Índice
    Producto Matriz-Matriz
    ¿Cómo explotar más la arquitectura?
    Buenas prácticas de programación CUDA
    OpenCL
    ¿Qué es OpenCL?
    ¿Qué diferencia OpenCL de CUDA?
    Ejemplo: Suma de vectores
    75

    Monografias.com

    ¿Qué es OpenCL?
    OpenCL: Open Computing Language

    Propuesto por a

    ¿Quién está involucrado?

    Monografias.com

    ¿Qué diferencia OpenCL de CUDA? (I)
    Punteros
    CUDA
    struct Node {Node* next}
    n=n->next
    Opencl
    struct Node {unsigned int next;}
    next=bufBase+n

    Monografias.com

    ¿Qué diferencia OpenCL de CUDA? (II)
    Kernels
    CUDA
    Programa compilado en formato binario
    OpenCL
    Se compila en tiempo de ejecución

    Palabras Clave y lenguaje utilizado para los kernel

    Monografias.com

    Ejemplo: Suma de vectores
    CUDA
    _global_ void
    SumaVec(const float *a, const float *b, float *c)
    // Índice al elemento del vector
    int indice=blockIdx.x*blockDim.x+threadIdx.x

    c[indice]=a[indice]+b[indice]
    }
    OpenCL
    _kernel void
    SumaVec(_global const float *a, _global const float *b, _global float *c)
    // Índice al elemento del vector
    int indice=get_global_id(0)

    c[indice]=a[indice]+b[indice]
    }

    Monografias.com

    Índice
    Producto Matriz-Matriz
    ¿Cómo explotar más la arquitectura?
    Buenas prácticas de programación CUDA
    OpenCL
    80

    Nota al lector: es posible que esta página no contenga todos los componentes del trabajo original (pies de página, avanzadas formulas matemáticas, esquemas o tablas complejas, etc.). Recuerde que para ver el trabajo en su versión original completa, puede descargarlo desde el menú superior.

    Todos los documentos disponibles en este sitio expresan los puntos de vista de sus respectivos autores y no de Monografias.com. El objetivo de Monografias.com es poner el conocimiento a disposición de toda su comunidad. Queda bajo la responsabilidad de cada lector el eventual uso que se le de a esta información. Asimismo, es obligatoria la cita del autor del contenido y de Monografias.com como fuentes de información.

    Categorias
    Newsletter