1. No category

Indice de contenidos
1. Introducción. [11 diapositivas]
2. Arquitectura. [25]
CUDA
1. El modelo hardware de CUDA. [6]
2. Las dos primeras generaciones (2006-2009) [7]
3. La tercera generación: Fermi (2010). [12]
Manuel Ujaldón Martínez
3. Programación. [20]
4. Sintaxis. [17]
5. Compilación. [4]
6. Depuración y optimización. [12]
7. Estrategias de mejora. [7]
8. Ejemplos. [3+12+8]
Departamento de Arquitectura de Computadores
Universidad de Málaga
El camino hacia los sombreadores
unificados que dieron lugar a CUDA
! Cuestión del año 2000: La programación de video-juegos
requiere instrumentos para recrear la riqueza de
materiales y efectos del mundo real.
! Respuesta del año 2002: Aparecen los sombreadores para
vértices y píxeles.
! Cuestión del año 2005: Algunas etapas del cauce de
segmentación gráfico son programables, pero con fuertes
restricciones:
I. Introducción
! En cómo pueden comunicarse entre ellas.
! En cómo pueden acceder a memoria global.
! Respuesta del año 2007: Se unifican los sombreadores, y
sobre ellos, surge el modelo de programación CUDA.
3
4
Alternativas a CUDA para
la programación paralela en GPUs
De menor a mayor afinidad durante esta década:
! Fujitsu’s: Llamadas a procedimientos remotos.
! Brook: El compilador pionero de Stanford.
! PeakStream: Librerías matemáticas para Radeon.
! RapidMind: Directivas de más alto nivel y una generación
de código más automatizada. Genera código para CPUs
x86, GPUs y Cell.
! Tilera’s: Hardware con topología malla escalable a 512
procesadores e incluso más.
! CAL (Compute Abstraction Layer) y CTM: Variantes AMD.
! OpenCL: El nuevo estándar, superconjunto de CUDA.
Los paradigmas de computación
hasta llegar a CUDA
5
¿Qué es CUDA?
“Compute Unified Device Architecture”
6
CUDA: Relación con la CPU
! CUDA ejecuta un programa sobre un device aparte (la
GPU), que actúa como coprocesador de un host (la CPU).
! Una plataforma diseñada conjuntamente a nivel software y
hardware para aprovechar la potencia de una GPU en
aplicaciones de propósito general.
! A nivel software, permite programar la GPU en C con mínimas
pero potentes extensiones SIMD para lograr una ejecución
eficiente y escalable.
! A nivel firmware, ofrece un driver para la programación
GPGPU que es compatible con el que se usa para renderizar.
! A nivel hardware, habilita múltiples niveles de paralelismo.
Según las 3 generaciones ya fabricadas, tenemos (1ª-2ª-3ª):
! Extensión de C.
! Biblioteca de funciones con tres tipos de componentes:
! Host: Control y acceso a los devices.
! Device: Funciones específicas de los devices.
! Comunes: Tipos vectoriales y un conjunto de rutinas soportadas por host y device.
CPU
HOST
! Multiprocesadores de 16-30-16 nodos sobre los que se ejecutan bloques.
! Nodos de 8-8-32 procesadores (SPs) sobre los que se ejecutan hilos.
! Jerarquía de memoria: Registros, memoria compartida y global (DRAM).
Memoria
principal
7
E/S
(PCIe)
GPU
DEVICE
Memoria de
vídeo
8
Ventajas de CUDA sobre la programación
GPGPU precursora (Cg y OpenGL)
Objetivos de CUDA
! Habilitar paralelismo masivo en GPU ejecutando
programas SIMD sin las limitaciones y sobrecargas del
API gráfico. GPGPU ya no es código OpenGL.
! Permitir al programador involucrarse a distintos grados
de exigencia según el rendimiento:
! Memoria direccionable a nivel de byte.
! Un hilo puede acceder a cualquier posición de memoria.
! Acceso ilimitado a memoria.
! Un hilo puede leer/escribir en tantas posiciones de memoria como
sea necesario.
! Básico: Posibilita una fácil portabilidad desde C++.
! Medio: Requiere un buen conocimiento de la arquitectura
gráfica.
! Avanzado: Permite un mapeo eficiente del problema sobre
muchos procesadores minimizando conflictos en el acceso a
memoria.
! Sincronización para los hilos y la memoria compartida
(para cada bloque).
! Los hilos pueden cargar datos en memoria compartida de forma
cooperativa dentro de un bloque.
9
47
10
Esfuerzo del programador CUDA
en relación al rendimiento esperado
Ventajas de CUDA sobre la programación
GPGPU precursora (2)
! La fase de aprendizaje resulta muy corta, ya que los
elementos a utilizar resultan familiares:
! Código: Apenas unas pocas extensiones al C convencional.
! Datos: No se requiere un conocimiento de gráficos (vértices,
texturas, píxeles).
! El tiempo de ejecución del código no sufre una
penalización por la sobrecarga del API gráfico (DirectX/
OpenGL).
! La depuración y optimización de código resultan mucho
más llevaderas.
48
11
Fuente: Curso de CUDA impartido en el MIT.
12
Si programamos con CUDA, la arquitectura
destino no tiene por qué ser CUDA
Si instalamos una arquitectura CUDA, la
programación tampoco tiene que ser CUDA
! A partir de CUDA 2.1 (2009), podemos compilar
indistintamente para many-core GPU o multi-core CPU
! Se soporta Windows, Linux y MacOS, y sobre ellos,
herramientas de desarrollo y lenguajes de programación
plenamente consolidados:
! Y con OpenCL, el abanico de plataformas es más amplio.
13
14
II.1. El modelo hardware
de CUDA
II. Arquitectura
15
16
El modelo hardware de CUDA:
Un conjunto de procesadores SIMD
Arquitectura del sistema de memoria
! La GPU consta de:
! Cada multiprocesador tiene:
GPU
! N multiprocesadores, cada uno
dotado de M procesadores.
Multiprocesador N
! Su banco de registros.
! Memoria compartida.
! Una caché de constantes y otra
de texturas, ambas de sólo
lectura y uso marginal.
Multiprocesador 2
Multiprocesador 1
! Paralelismo masivo:
Procesador 1
! Aplicado sobre miles de hilos.
! Compartiendo datos a
diferentes niveles.
Procesador 2
…
G80
Período de
vigencia
! Computación heterogénea,
complementa CPU y GPU:
! GPU: Intensiva en datos.
Paralelismo de grano fino.
! CPU: Gestión y control.
Paralelismo de grano grueso.
ProcesadorM
Unidad de
Control
2006-07 2008-09 2010-11
N
16
30
16
M
8
8
32
128
240
512
Número total
de cores
! La memoria global es la
memoria de vídeo (DRAM):
GTX200 Fermi
17
! Tres veces más rápida que la
de la CPU, pero...
! ... ¡500 veces más lenta que la
memoria compartida! (que es
SRAM en realidad).
GPU
Multiprocesador N
Multiprocesador 2
Multiprocesador 1
Memoria compartida
Registros
Registros
Procesador 1
Procesador 2
Registros
…
Procesador M
Unidad
de
Control
Caché para
constantes
Caché para
texturas
Memoria global
Latencia y ancho de banda
de la memoria en CUDA
Arquitectura de cada grupo o nodo básico
para la construcción de las distintas GPUs
! Memoria de la CPU
! Se compone de un par de multiprocesadores de ocho
cores o SPs (stream processors). Cada multiprocesador
dispone de una memoria compartida o caché de 16 KB
casi tan rápida como el banco de registros, y por la que
pueden comunicarse los hilos pertenecientes al mismo.
! Computación de punto flotante:
! Ancho de banda con mem. vídeo: 3.2 GB/s.(PCIe) y 5.2 GB/s(PCIe2).
! Memoria de vídeo global/local
! Gran ancho de banda (80-100 GB/s) y latencia, no pasa por caché.
! Memoria compartida
! Baja latencia, ancho de banda muy elevado, tamaño reducido.
! Actúa como una caché gestionada por el usuario (scratchpad).
! Memoria de texturas/constantes
! Para 32 bits (2006).
! Para 64 bits (2008).
! De sólo lectura, alta/baja latencia, pasa por caché, uso marginal.
Host
Device
CPU
Chipset
DRAM
Local
Global
Constantes
DRAM
Texturas
13
18
GPU
Multiprocesador
Registros y
memoria compartida
Cachés para constantes
y texturas
19
20
Escalabilidad del modelo
aumentando el número de nodos
Escalabilidad del modelo: En 8 y 15 nodos
llegamos a las dos primeras generaciones
Dos nodos: GeForce 8400
! El modelo de 8 nodos es el de la GPU G80 (1ª gener.).
Cuatro nodos: GeForce 8600
! El modelo de 15 nodos es el de la GTX200 (2ª gener.).
21
22
La arquitectura en general
II.2. Las dos primeras
generaciones
! Se compone de multiprocesadores de 8 cores, donde los
GFLOPS escalan con el número de cores (stream processors),
y el ancho de banda escala con el número de controladores
de memoria según el modelo comercial:
32, 128 ó 240 cores
Gestor de hilos
GDDR3
>100 GB/sg.
Número de
multiprocesadores
512 bit
Memoria
de vídeo
23
8600 9800 GTX
GTS GTX 200
4
16
30
93
429
624
Controladores de
memoria
2
4
8
Ancho de banda
(GB/s)
32
70
141
GFLOPS
24
Detalle de cada multiprocesador:
Sus unidades funcionales
! Ocho ALUs (una en cada
procesador escalar) para
operaciones add/mul, tanto
para enteros como para p.f.
! Dos SFUs para operadores
complejos (sin, cos, inv, sqrt).
! Una unidad p.f. doble precisión.
! Rendimiento pico en cada multiprocesador:
Inestabilidades numéricas
(ya han sido resueltas en Fermi)
Búsqueda de instrucción
Planif. instr.
Format
Registros
Rounding modes for
FADD and FMUL
Memoria compartida
ALU
Denormal handling
ALU
ALU
NaN support
ALU
ALU
Overflow and
infinity support
ALU
SFU
ALU
SFU
Unidad de punto flotante
en doble precisión
ALU
Flags
! En la G80: 18 ops. por ciclo (8 mul/add flotante y 2 ops complejas):
! Para 16 multiprocesadores a 1.35 GHz: 16 x 18 x 1.35 = 388.8 GFLOPS
! En la GTX200: 24 ops. por ciclo (1 mul/add flotante y 1 mul):
! Para 30 multiprocesadores a 1.296 GHz: 30 x 24 x 1.296 = 933 GFLOPS
25
Detalle de cada procesador o SP
G80
SSE
IBM Altivec
Cell SPE
IEEE 754
IEEE 754
IEEE 754
IEEE 754
Round to nearest and
round to zero
All 4 IEEE, round to
nearest, zero, inf, -inf
Round to nearest only
Round to zero/truncate
only
Flush to zero
Supported,
1000’s of cycles
Supported,
1000’s of cycles
Flush to zero
Yes
Yes
Yes
No
Yes, only clamps to
max norm
Yes
Yes
No, infinity
No
Yes
Yes
Some
Square root
Software only
Hardware
Software only
Software only
Division
Software only
Hardware
Software only
Software only
Reciprocal estimate
accuracy
24 bit
12 bit
12 bit
12 bit
Reciprocal sqrt
estimate accuracy
23 bit
12 bit
12 bit
12 bit
log2(x) and 2^x
estimates accuracy
23 bit
No
12 bit
No
26
La primera generación: G80 (GeForce 8800)
! Se trata de unidades vectoriales multihilo:
GPU G80 (en torno a 600 MHz, frecuencia muy inferior a la de sus cores)
Multiprocesador 16
La emisión múltiple de
instrucciones se consigue
ejecutando más veloz el
front-end (CPI=2) que el
back-end (CPI = 4,16,32).
Multiprocesador 2
Multiprocesador 1 (los bloques de código CUDA se mapean sobre los multipr.)
Memoria compartida (16 KB)
Registros
Core 1
Esta idea recuerda al
hyperthreading de Intel.
(1.35 GHz)
Registros
Core 2
…
Caché de texturas
Registros
Core 8
Unidad de
control
(emite
instrucciones
SIMD)
(los kernels se mapean
sobre los cores)
Memoria global (hasta 1.5 GB) (GDDR3 @ 2x 800MHz)
27
28
Escalabilidad para futuras generaciones:
Alternativas para su crecimiento futuro
La segunda generación (GTX200)
GPU GTX200 (en torno a 600 MHz)
Multiprocesador 30
Multiprocesador 2
Multiprocesador 1 (los bloques de código CUDA se mapean sobre los multipr.)
Memoria compartida (16 KB)
Registros
Core 1
(1.30 GHz)
Registros
Core 2
…
Caché de texturas
Registros
Core 8
Unidad de
control
(emite
instrucciones
SIMD)
(los kernels se mapean
sobre los cores)
Memoria global (hasta 4 GB) (GDDR3, 512 bits @ 2x 1.1GHz = 141.7 GB/s)
29
GPU
! Aumentar el número de
Multiprocesador 30
multiprocesadores por pares
(escalabilidad en 2ª gener.)
(nodo básico), esto es, crecer
Multiprocesador 2
en la dimensión Z. Es lo que
Multiprocesador 1
hizo la 2ª gener. (de 16 a 30).
Memoria compartida
! Aumentar el número de
Registros
Registros
Registros
procesadores de cada
Core 2 … Core 8
Core 1
multiprocesador, o crecer en la
(escalabilidad en 3ª gener.)
dimensión X. Es lo que hizo la
Caché de texturas
3ª gener. (de 8 a 32).
! Aumentar el tamaño de la
memoria compartida, esto es,
Memoria global
crecer en la dimensión Y.
30
La nueva GPU en cifras, y respecto
a sus precedesores en la saga GeForce
II. 3. La tercera generación:
Fermi
G80 (2006)
GTX200 (2008)
Fermi (2010)
Transistores
681 millones
1400 millones
3000 millones
Cores CUDA
128
240
512
Vel. cálculo en
doble precisión
Ninguna
30 madds/clock
256 madds/clock
Planificadores
de warps
1
1
2
Shared memory
31
16 KB
16 KB
Caché L1
Ninguna
Ninguna
16 KB + 48 KB
(o vice versa)
Caché L2
Ninguna
Ninguna
768 KB
Corrección de
errores (DRAM)
No
No
Sí
Anchura del bus
de direcciones
32 bits
32 bits
64 bits
32
DRAM
I/F
!
- 3000 millones de transistores.
!
- Más del doble de cores que en la
generación anterior (512 vs. 240).
!
- 8 veces más de rendimiento pico
sobre punto flotante de doble
precisión que la generación anterior.
A 2 GHz, proporciona 1 TFLOPS.
!
- Corrección de errores (ECC).
!
- Cachés L1 y L2 transparentes.
!
- El doble de ancho de banda con
memoria de vídeo (GDDR5).
!
- Hasta 1 Terabyte de memoria de
vídeo.
!
- Kernels concurrentes, C++.
L2
DRAM
I/F
DRAM
I/F
DRAM Giga ThreadHOST I/
I/F
DRAM
I/F
DRAM
I/F
Principales rasgos de Fermi
“The soul of a Supercomputer in the body of a GPU”
La disposición física de los
cores y la memoria apenas
ha cambiado respecto a su
predecesor GTX200.
Arquitectura global de Fermi
33
34
Arquitectura de los multiprocesadores
! Mantener la escalabilidad de la GPU en años venideros.
! Atraer a un mayor número de usuarios y aplicaciones.
! Se dispone de 16 multiprocesadores,
cada uno compuesto de 32 cores.
! Gran mejora de la computación de
punto flotante en doble precisión:
L2
! Aumenta en un factor 8x su
rendimiento pico vs. G80.
! Tan sólo computa el doble de
lento que en simple precisión.
DRAM I/F
DRAM I/F
DRAM I/F
DRAM I/F Giga Thread HOST I/F
DRAM I/F
DRAM I/F
Objetivo de Fermi
! Dos planificadores de threads.
! 64 KB. de SRAM que se reparten entre
la memoria compartida y la caché L1.
35
Load/Store Units x 16
Special Func Units x 4
36
Arquitectura de los cores
La jerarquía de memoria
! Implementa el formato IEEE-754 en su
versión de 2008, aventajando incluso
a las CPUs más avezadas.
! Instrucción “madd” (suma
y producto simultáneos)
disponible tanto para simple
como para doble precisión.
FP Unit
! ALU rediseñada para optimizar
las operaciones sobre enteros
de 64 bits y admitir operaciones
de precisión extendida.
! La primera GPU que ofrece
una caché L1 típica on-chip,
que combina con la shared
memory de CUDA para un
total de 64 KB por cada
multiprocesador (32 cores).
! También incluye una caché
unificada de 768 KB con
coherencia de datos para el
conjunto de cores.
INT Unit
Load/Store Units x 16
Special Func Units x 4
13
37
38
Planificador hardware de threads GigaThread
(HTS)
El planificador de instrucción
! Gestiona jerárquicamente miles de threads
simultáneamente activos.
! El cambio de contexto de la aplicación es diez veces más
rápido.
HTS
14
39
40
Planificador hardware de threads GigaThread
(HTS)
GigaThread: El motor de transferencia
de datos streaming GigaThread (SDT)
! Dos coprocesadores de DMA:
! Gestiona jerárquicamente miles de
threads simultáneamente activos.
SDT
! Ejemplo de actividad:
!
Kernel 0
SDT0
Kernel 1
SDT1
SDT0
Kernel 2
SDT1
SDT0
Kernel 3
SDT1
SDT0
SDT1
41
42
El entorno de desarrollo Nexus: Integrando
toda la programación en MS Visual Studio
III. Programación
15
43
44
El ecosistema corporativo y funcional
de CUDA en su vertiente software
El modelo de programación CUDA
! La GPU (device) ofrece a la CPU (host) la visión de un
coprocesador altamente ramificado en hilos.
! Que tiene su propia memoria DRAM.
! Donde los hilos se ejecutan en paralelo sobre los núcleos (cores o
stream processors) de un multiprocesador.
GPU
Multiprocesador 1
Multiprocesador 2
Multiprocesador N
! Los hilos de CUDA son extremadamente ligeros.
Oil & Gas
Finance
CFD
Medical
Biophysics
Imaging
Numerics
DSP
! Se crean en un tiempo muy efímero.
! La conmutación de contexto es inmediata.
ANEO
! Objetivo del programador: Declarar miles de hilos, que
la GPU necesita para lograr rendimiento y escalabilidad.
GPU Tech
EDA
45
46
!
Ejemplo: GeForce 8 y Quadro FX
Estructura de un programa CUDA
! El número de procesadores, su velocidad y la memoria
global disponibles depende del modelo comercial que
suscribamos.
! Cada multiprocesador procesa lotes de bloques, uno
detrás de otro
Número de
multiprocesadores
Frecuencia de
reloj (GHz)
Memoria de vídeo
(MB)
GeForce 8800 GTX
16
1.35
768
GeForce 8800 GTS
12
1.2
640
Quadro FX 5600
16
1.35
1500
Quadro FX 4600
12
1.2
768
Modelo de GPU
! Bloques activos = los bloques procesados por un multiprocesador
en un lote.
! Hilos activos = todos los que provienen de los bloques que se
encuentren activos.
! Los registros y la memoria compartida de un multiprocesador se reparten entre sus hilos activos. Para un
kernel dado, el número de bloques activos depende de:
! El número de registros requeridos por el kernel.
! La memoria compartida consumida por el kernel.
47
48
Recursos y limitaciones según la GPU
que utilicemos para programar CUDA
Conceptos básicos
!
!
!
!
!
Los programadores se enfrentan al reto de exponer el
paralelismo para múltiples cores y para múltiples hilos por
core. Para ello, deben usar los siguientes elementos:
Device = GPU = conjunto de multiprocesadores.
Multiprocesador = conjunto de procesadores y memoria
compartida.
Kernel = programa ejecutándose en GPU.
Grid = matriz de bloques de hilos que ejecutan un kernel.
Bloque de hilos (thread block) = grupo de hilos SIMD que
ejecutan un kernel delimitando su dominio computacional
según su threadID, y que pueden comunicarse a través de la
memoria compartida del multiprocesador.
Parámetro
CUDA Compute Capabilities
Máximo
8 bloques
!
····
Máximo
512 hilos
····
· · · · ·· ·· · ·
····
Asignación
a un multiproc.
HW.
Escalabilidad
8
8
32
HW.
Escalabilidad
32
32
32
SW.
Throughput
8
8
8
SW.
Throughput
Hilos / Bloque
512
512
512
SW.
Paralelismo
Hilos / Multiprocesador
768
1 024 1 536
SW.
Paralelismo
Registros de 32 bits / Multiprocesador
8 192
16 384 4 096
HW.
Working set
Memoria compartida / Multiprocesador
16 384
16K
48K
HW.
Working set
SFU
Core
Core
Core
Core
SFU
Máximo 768, 1024 ó 1536
hilos según la gener. de GPU
! Los hilos se asignan a los multiprocesadores en “bloques”,
que constituyen la unidad de asignación de hilos.
! Cada multiprocesador puede tener hasta 8 bloques y cada
bloque hasta 512 hilos. En total, un máximo de 768, 1024 o
1536 hilos pueden asignarse a cada multiprocesador.
! Los hilos de un bloque comparten información a través de
memoria compartida, y se sincronizan mediante barreras.
CPU (host)
...
32 hilos
=
Multiproc.
32 hilos
32 hilos
Bloque de hilos
51
GPU (device)
Grid 1
Bloque
(0, 0)
Bloque
(1, 0)
Bloque
(2, 0)
Bloque
(0, 1)
Bloque
(1, 1)
Bloque
(2, 1)
Grid 2
Kernel 2
Bloque (1, 0)
Hilo
(0, 0)
Hilo
(0, 1)
Hilo
(0, 2)
Warps
16
50
Kernel 1
! Cada bloque activo se descompone
en WARPs o grupos de 32 hilos con
ID correlativo que se ejecutan
usando tiempo compartido en un
multiprocesador.
! Los hilos de un WARP se
ejecutan físicamente en paralelo.
! Los WARPs y los bloques se
ejecutan lógicamente en paralelo.
Core
Core
16 384
49
WARPs. Concepto
Core
Core
16
Bloques de hilos / Multiprocesador
Memoria compartida
Impacto
30
Hilos / Warp
Registros
Limitación
16
Procesadores / Multiprocesador
B1
!!
1.0 y 1.1 1.2 y 1.3 Fermi
Multiprocesadores / GPU
Planificación de instrucciones:
Bloques de hilos
Bn
Valor según gener. GPU
Warp 0
…
Warp 2
…
Warp 4
…
Hilo
(31, 0)
Hilo
(32, 0)
Hilo
(31,
Hilo
(32, 1)
Hilo
(31, 2)
Hilo
(32, 2)
Warp 1
…
Warp 3
…
Warp 5
…
Hilo
(63, 0)
Hilo
(63,
Hilo
(63, 2)
DRAM
16
Half Warps
52
Device
Memory
WARPs. Planificación
WARPs. Ejecución
Multiprocesador
! Dado que el tamaño del WARP es de 32 hilos:
! Si hay 8 procesadores en cada multiprocesador (como en las series
8, 9 y 10 de GeForce), cada WARP de una instrucción consume 4
ciclos de reloj en su ejecución.
! Si hay 32 procesadores en cada multiprocesador (a partir de Fermi),
cada WARP de una instrucción consume un solo ciclo.
Warp j (32 hilos)
Bi
Búsqueda de instrucción
L1 datos
Core
Core
SFU
SFU
Warp 5
Instr. 11
Warp 12
Instr. 3
Warp 12
Instr. 4
Warp 11
Instr. 8
Warp 15
Instr. 1
ciclos
! Round-robin/aging para seleccionar el próximo WARP a planificar de
entre aquellos con operandos ya leídos.
! Scoreboarding para evitar riesgos en el análisis de dependencias.
Registros
Memoria compartida
Core
Warp 23
Instr. 17
! El WARP es la unidad de planificación. Se usa:
Instrucción x de warp j
····
······
····
Warp 5
Instr. 10
Core
Core
Core
Core
Core
! El cambio de contexto entre WARPs de un multiprocesador
se lleva a cabo sin penalidad en ciclos de ejecución.
53
Recopilando sobre kernels,
bloques y paralelismo
! Los kernels se lanzan en grids.
! Sólo se ejecuta un kernel a la vez.
! Un bloque se ejecuta en un
multiprocesador.
Escalabilidad transparente
! Dado que los bloques de hilos no pueden sincronizarse,
el hardware tiene libertad para acomodar la ejecución de
un bloque de hilos en cualquier multiprocesador en
cualquier instante.
Grid
Bloque (0, 0)
Bloque (1, 0)
! El bloque no migra.
! Varios bloques pueden residir
concurrentemente en un
multiprocesador.
! Con las consabidas limitaciones:
! 8 bloques concurrentes.
! 512 hilos en cada multiprocesador.
! 768/1024/1536 hilos en total.
Memoria compartida
Memoria compartida
Kernel grid
Device
Regs
Hilo (0, 0)
Regs
Hilo (1, 0)
Regs
Hilo (0, 0)
Regs
Block 0
Hilo (1, 0)
! Y otras limitaciones entran en juego
debido al uso conjunto de la memoria:
! El banco de registros se particiona
entre todos los hilos existentes.
! La memoria compartida se particiona
entre todos los bloques de hilos.
54
Block 2
Memoria global
55
Block 1
Block 3
Block 4
Block 5
Block 6
Block 7
Device
Block 0
Block 1
Block 2
Block 3
Block 4
Block 5
Block 6
Block 7
Block 0
Block 1
Block 2
Block 3
Block 4
Block 5
Block 6
Block 7
! Un kernel puede ejecutarse
sobre cualquier número de multiprocesadores.
! El lanzamiento de un kernel sirve como punto
de sincronización para los bloques.
56
Kernels (y su relación con los hilos)
Hilos (y su relación con los bloques)
! Las porciones paralelas de una aplicación que corre en
la CPU se ejecutan en la GPU como kernels.
! Sólo un kernel se ejecuta en un momento dado en una
GPU (esto es, NO hay multiproceso de kernels).
! Cuando el kernel finaliza, todos los recursos de la GPU
se liberan y quedan disponibles íntegramente para el
kernel siguiente.
! La cooperación entre hilos resulta muy valiosa:
! Según su threadID, cada hilo:
! Ejecuta el mismo código sobre un
área diferente de datos.
! Puede tomar decisiones de control
para diferenciar su ejecución del resto.
0 1 2 3 4 5 6 7
! Comparten resultados para ahorrar computaciones.
! Comparten accesos a memoria de vídeo para reducir drásticamente
el ancho de banda (y el consumo del chip).
! El bloque garantiza rendimiento y escalabilidad, ya que
permite replicar la ejecución de un grupo de hilos tantas
veces como sea necesario en función del volumen de datos:
! Permitiendo mantener el paralelismo de grano fino.
! Sin penalidad, ya que el cambio de contexto es gratis.
threadID
…
float x = input[threadID];
float y = func(x);
output[threadID] = y;
…
Thread Block 0
threadID
0
1
2
3
4
5
Thread Block 0
6
7
…
float x = input[threadID];
float y = func(x);
output[threadID] = y;
…
57
0
1
2
3
4
5
Thread Block N-1
6
7
…
float x = input[threadID];
float y = func(x);
output[threadID] = y;
…
0
…
1
2
3
4
5
6
7
…
float x = input[threadID];
float y = func(x);
output[threadID] = y;
…
Particionamiento de computaciones y datos
Manipulación de datos
! Un bloque de hilos es un lote
de hilos que pueden cooperar:
! Constituye una de las diferencias más importantes entre la
CPU y la GPU, y una de las principales razones para el mayor
rendimiento pico que atesora la GPU.
! El programador gestiona de forma explícita la caché.
! La caché es un millar de veces más pequeña en la GPU, por
lo que el programador debe explotar al máximo la localidad.
! Compartiendo datos a través de
memoria compartida.
! Sincronizando su ejecución para
acceder a memoria sin conflictos.
! Un kernel se ejecuta como una
malla o grid 1D o 2D de
bloques de hebras 1D, 2D o 3D.
! Los hilos y los bloques tienen
IDs para que cada hilo pueda
acotar sobre qué datos trabaja,
y simplificar el dir. a memoria al
procesar datos multidim.
CPU (host)
GPU (device)
Grid 1
Kernel 1
Bloque
(0, 0)
Bloque
(1, 0)
Bloque
(2, 0)
Bloque
(0, 1)
Bloque
(1, 1)
Bloque
(2, 1)
Grid 2
Kernel 2
Memoria Ubicación Caché
Bloque (1, 1)
Hilo
Hilo
(0, 0) (1, 0)
Hilo
(2, 0)
Hilo
Hilo
(3, 0) (4, 0)
Hilo
Hilo
Hilo
Hilo
Hilo
(0, 1) (1, 1) (2, 1) (3, 1) (4, 1)
Hilo
Hilo
Hilo
Hilo
Hilo
(0, 2) (1, 2) (2, 2) (3, 2) (4, 2)
59
Acceso
Ámbito
Lect./escr.
Un hilo
58
Declaración Vigencia
Local
Off-chip
No
__device__
Hilo
Compartida
On-chip
-
Lect./escr. Hilos de un bloque
__shared__
Bloque
Global
Off-chip
No
Lect./escr. Los hilos y la CPU
__global__
Kernel
Constantes
Off-chip
Sí
Lectura
Los hilos y la CPU __constant__
Kernel
De texturas
Off-chip
Sí
Lectura
Los hilos y la CPU
Kernel
__texture__
60
Tipos de memoria y su jurisdicción
Espacios de memoria (1)
Registros
Hilo
Bloque
! La CPU y la GPU tiene espacios de memoria separados:
Integrada
en la GPU
Memoria
compartida
!!!
! Para comunicar ambos procesadores, se utiliza el bus PCI-express.
! En la GPU se utilizan funciones para alojar memoria y copiar datos
de la CPU de forma similar a como la CPU procede en lenguaje C.
! Los hilos pueden compartir el espacio de memoria
compartida para comunicarse entre ellos.
Grid 0
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
!!!
! Los punteros son sólo direcciones:
Memoria
Memoria
para
global
(lectura y constantes
(lectura)
escritura)
Grid 1
Memoria
para
texturas
(lectura)
! No se puede conocer a través del valor de un puntero si la dirección
pertenece al espacio de la CPU o al de la GPU.
! Hay que ir con mucha cautela a la hora de acceder a los datos a
través de punteros, ya que si un dato de la CPU trata de accederse
desde la GPU o viceversa, el programa se colgará.
Externa
a la GPU
(pero
dentro
de la
tarjeta
gráfica)
! El nivel de memoria global es el único visible a todos los entes
definidos por el programador (hilos, bloques, kernels y grids).
61
Principales debilidades/riesgos de CUDA
Espacios de memoria (2)
Grid en GPU
! Cada hilo puede:
! Leer/escribir registros per-thread
! Leer/escribir mem. compartida per-block
! Leer/escribir mem. global per-grid
! Cada hilo también puede, por motivos de
comodidad o rendimiento:
! Leer/escribir memoria local per-thread
! Leer memoria de constantes per-grid
! Leer memoria de texturas per-grid
! La CPU puede:
! Leer/escribir en memoria
global, constante y de texturas
(mapeadas sobre DRAM)
Bloque (0, 0)
Memoria compartida
Memoria compartida
Registros
Registros
Registros
Registros
Hilo (0, 0)
Hilo (1, 0)
Hilo (0, 0)
Hilo (1, 0)
Memoria
local
Memoria
local
! El ancho de banda entre memoria global (la de la tarjeta
gráfica) y los procesadores puede saturarse fácilmente.
Las tareas que tienen un bajo índice de reutilización de
datos se quedan hambrientas. Ej: Suma vector-escalar,
producto de vectores.
! Limitada capacidad del banco de registros y la memoria
compartida que comparten todos los hilos de un
multiprocesador. Utilizar el CUDA Occupancy Calculator
para ayudarse en la toma de decisiones.
! Los saltos condicionales degradan notablemente el
rendimiento si no se estructuran de forma sabia.
Bloque (1, 0)
Memoria
local
Memoria
local
Memoria global
CPU
(host)
62
Memoria de constantes
Memoria de texturas
63
64
Cinco claves para maximizar
el rendimiento del código
1. Expresar explícitamente todo el paralelismo posible aplicando
SIMD de grano fino para definir multitud de hilos. Recordar
que el cambio de contexto es gratis en CUDA.
1. Si los hilos de un mismo bloque necesitan comunicarse,
utilizar la memoria compartida y __syncthreads()
2. Si los hilos de diferentes bloques necesitan comunicarse,
utilizar la memoria global y descomponer la computación
en múltiples kernels.
2. Aprovechar el ancho de banda con memoria: Pocas
transferencias grandes en lugar de muchas pequeñas.
3. Optimizar la localidad de acceso: Reutilización de datos.
4. Ocultar latencias con memoria global maximizando la
ocupación de unidades funcionales. Intensidad aritmética.
5. Maximizar el CPI del código (throughput): Seleccionar la
instrucción de menor latencia en el repertorio CUDA.
IV. Sintaxis
65
CUDA es C con algunas palabras clave más.
Un ejemplo preliminar
66
La interacción entre la CPU y la GPU
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];
Código C estándar
}
// Invoke serial SAXPY kernel
saxpy_serial(n, 2.0, x, y);
Código CUDA equivalente de ejecución paralela en GPU:
__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];
}
// Invoke parallel SAXPY kernel with 256 threads/block
int nblocks = (n + 255) / 256;
saxpy_parallel<<<nblocks, 256>>>(n, 2.0, x, y);
67
! CUDA extiende el lenguaje C con un nuevo tipo de función,
kernel, que ejecutan en paralelo los hilos activos en GPU.
! El resto del código es C nativo que se ejecuta sobre la CPU
de forma convencional.
! De esta manera, el típico main() de C combina la ejecución secuencial en CPU y paralela en GPU de kernels CUDA.
! Un kernel se lanza siempre de forma asíncrona, esto es, el
control regresa de forma inmediata a la CPU.
! Cada kernel GPU tiene una barrera implícita a su conclusión,
esto es, no finaliza hasta que no lo hagan todos sus hilos.
! Aprovecharemos al máximo el biprocesador CPU-GPU si les
vamos intercalando código con similar carga computacional.
68
Identificación de los entes al programar
(sobre un ejemplo de suma de matrices)
La interacción entre la CPU y la GPU (cont.)
! BlockIdx: vector (1D o 2D) que identifica el
bloque dentro del grid.
CPU
CPU
! ThreadIdx: vector (1D, 2D o 3D) que identifica
el hilo dentro de su bloque.
GPU
CPU
__global__ void matAdd (float A[N][N],float B[N][N],float C[N][N])
{
int i = blockIdx.x*blockDim.x + threadIdx.x;
int j = blockIdx.y*blockDim.y + threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
{
int main(){
dim3 dimBlock(4,4);
dim3 dimGrid (N/dimBlock.x, N/dimBlock.y);
matAdd <<< dimGrid, dimBlock >>> (A, B, C);
69
{
{
blockIdx = (3,0)
threadIdx = (3,2)
BlockDim.x
BlockIdx.y
es 0
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
BlockIdx.x
es 3
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
Grid de
bloques
}
70
! Modificadores para las funciones ejecutadas en la GPU:
! Cada multiprocesador ejecuta los 32 hilos de un warp
según el paradigma SIMD, esto es, todos los hilos ejecutan
la misma instrucción.
! Si las bifurcaciones afectan de forma diferente a los hilos
de un warp (divergencia), se secuencializa la ejecución de
los hilos del warp (es SIMD, no SPMD), con la consiguiente
pérdida de rendimiento.
! Para evitar esto, existen algunos trucos:
if (ThreadIdx / WARP_SIZE >2)
¿Qué hilo computa C[15,2]?
Modificadores para las funciones y
lanzamiento de ejecuciones en GPU
Control de flujo para el programa
if (ThreadIdx.x > 2)
Las tres matrices se
particionan igual,
otorgando un
elemento a cada hilo
! Para cada uno de los 4x4 hilos del bloque:
GPU
Un kernel no comienza su ejecución en GPU hasta que
no hayan finalizado todas las llamadas CUDA anteriores.
}
Matrices A, B y C
de 16x16 elementos
! Para cada uno de los 4x4 bloques del grid:
blockDim.y
Ejecución
__global__ kernelA(){···}
__global__ kernelB(){···}
int main()
···
kernelA <<< dimGridA, dimBlockA >>> (params.);
···
kernelB <<< dimGridB, dimBlockB >>> (params.);
···
! __global__ void MyKernel() { } // Invocado por la CPU
! __device__ float MyFunc() { } // Invocado por la GPU
! Modificadores para las variables que residen en la GPU:
! __shared__ float MySharedArray[32]; // En mem. caché
! __constant__ float MyConstantArray[32];
! Configuración de la ejecución para lanzar kernels:
! dim2 gridDim(100,50);
// 5000 bloques de hilos
! dim3 blockDim(4,8,8);
// 256 bloques de hilos
! MyKernel <<< gridDim,blockDim >>> (pars.); // Lanzam.
}
Con divergencia
Sin divergencia
71
72
Para conocer en tiempo de ejecución
los recursos hardware con los que contamos
Variables intrínsecas y funciones accesibles
! dim3 gridDim;
! dim3 blockDim;
! Cada GPU disponible en la capa hardware recibe un número
entero consecutivo que la identifica, comenzando por el 0.
! Para conocer el número de GPUs disponibles:
// Dimensión del grid
// Dimensión del bloque
! cudaGetDeviceCount(int* count);
! Para conocer los recursos disponibles en la GPU dev
(caché, registros, frecuencia de reloj, ...):
! uint3 blockIdx; // Indice del bloque dentro de la malla
! uint3 threadIdx; // Indice del hilo dentro del bloque
! cudaGetDeviceProperties(struct cudaDeviceProp* prop, int dev);
! Para conocer la mejor GPU que reúne ciertos requisitos:
! void __syncthreads(); // Sincronización entre threads
! cudaChooseDevice(int* dev, const struct cudaDeviceProp* prop);
! Para seleccionar una GPU concreta:
El programador debe elegir el tamaño del bloque
y el número de bloques para explotar al máximo
el paralelismo del código durante su ejecución.
! cudaSetDevice(int dev);
! Para conocer en qué GPU estamos ejecutando el código:
73
Sincronización entre la CPU y la GPU
cuando disponemos de varias GPUs
! cudaGetDevice(int* dev);
74
Para gestionar la memoria de vídeo
! Cada hilo de la CPU sólo puede controlar una GPU, por lo
tanto, necesitamos declarar tantos hilos en CPU (usando
pthreads) como GPUs queramos poner a trabajar.
! Una CPU no puede controlar varias GPUs, pero una GPU sí
puede ser accedida desde varios hilos de la CPU. El driver
de la GPU se encarga de serializar las llamadas.
! Los recursos CUDA alojados desde un hilo de la CPU sólo
pueden ser utilizados por llamadas CUDA procedentes de
ese mismo hilo. Ejemplo de mal uso:
! Para reservar y liberar memoria en la GPU:
y cudaFree(p)
! Para mover áreas de memoria entre CPU y GPU,
tras declarar malloc(h_A) en la CPU y
cudaMalloc(d_A) en la GPU:
! cudaMalloc(void* p, size_t numBytes)
! Desde la CPU a la GPU:
! cudaMemcpy(d_A, h_A, numBytes, cudaMemcpyHostToDevice);
! Desde la GPU a la CPU:
! cudaMemcpy(h_A, d_A, numBytes, cudaMemcpyDeviceToHost);
! El hilo 2 de la CPU aloja memoria en GPU y almacena su dirección
en el puntero P.
! El hilo 3 realiza una llamada a CUDA pasando P como parámetro.
75
76
Ejemplo 1: Descripción
Ejemplo 1: Implementación
int main()
{
int dimx = 16;
int num_bytes = dimx*sizeof(int);
int *d_a=0, *h_a=0; // device and host pointers
! Alojar n enteros en la memoria de la CPU.
! Alojar n enteros en la memoria de la GPU.
! Inicializar la memoria de la GPU a cero.
! Copiar los valores desde la GPU a la CPU.
! Imprimir los valores.
h_a = (int*)malloc(num_bytes);
cudaMalloc( (void**)&d_a, num_bytes );
if( 0==h_a || 0==d_a ) printf("couldn't allocate memory\n");
cudaMemset( d_a, 0, num_bytes );
cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
for(int i=0; i<dimx; i++) printf("%d ", h_a[i] );
free( h_a );
cudaFree( d_a );
77
}
78
Ejemplo 2: Incrementar un valor “b”
a los N elementos de un vector
Transferencias de memoria asíncronas
! Las llamadas a cudaMemcpy() son síncronas, esto es:
Programa C en CPU
! No comienzan hasta que no hayan finalizado todas las llamadas
CUDA que le preceden.
! El retorno a la CPU no tiene lugar hasta que no se haya realizado la
copia en memoria.
void increment_cpu(float *a, float b, int N)
{
for (int idx = 0; idx<N; idx++)
a[idx] = a[idx] + b;
! A partir de CUDA Compute Capabilities 1.2 es posible
utilizar la variante cudaMemcpyAsync(), cuyas
diferencias son las siguientes:
}
void main()
{
.....
increment_cpu(a, b, N);
}
! El retorno a la CPU tiene lugar de forma inmediata.
! Podemos solapar comunicación y computación.
! En la sección “Estrategias de mejora” pondremos un ejemplo.
79
Programa CUDA en GPU
__global__ void increment_gpu(float *a, float b, int N)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N)
a[idx] = a[idx] + b;
}
void main()
{
…..
dim3 dimBlock (blocksize);
dim3 dimGrid( ceil( N / (float)blocksize) );
increment_gpu<<<dimGrid, dimBlock>>>(a, b, N);
}
62
80
Código en CPU para el ejemplo 2
(azul es C, verde es CUDA, rojo son vars.)
Ejemplo 2: Incrementar un valor “b”
a los N elementos de un vector
Extensiones
al lenguaje
Con N=16 y blockDim=4, tenemos 4 bloques de hilos,
encargándose cada hilo de computar un elemento del vector.
blockIdx.x=0
blockDim.x=4
threadIdx.x=0,1,2,3
idx=0,1,2,3
blockIdx.x=1
blockDim.x=4
threadIdx.x=0,1,2,3
idx=4,5,6,7
blockIdx.x=2
blockDim.x=4
threadIdx.x=0,1,2,3
idx=8,9,10,11
int idx = (blockId.x * blockDim.x) + threadIdx.x;
Se mapeará del índice local threadIdx al índice global
// aloja memoria en la GPU
float* d_A = 0; cudaMalloc((void**)&d_A, numbytes);
// copia los datos de la CPU a la GPU
cudaMemcpy(d_A, h_A, numBytes, cudaMemcpyHostToDevice);
blockIdx.x=3
blockDim.x=4
threadIdx.x=0,1,2,3
idx=12,13,14,15
// ejecuta el kernel. B es es número de bloques
increment_gpu <<< N/blockSize, blockSize >>> (d_A, b);
// copia los datos de regreso a la CPU
cudaMemcpy(h_A, d_A, numBytes, cudaMemcpyDeviceToHost);
Patrón de acceso
común
Nota: blockDim debería ser >= 32 (warp size) en código real, esto es sólo un ejemplo
63
// aloja memoria en la CPU
unsigned int numBytes = N * sizeof(float);
float* h_A = (float*) malloc(numBytes);
81
// libera la memoria de vídeo
cudaFree(d_A);
64
Interoperabilidad con los programadores
de los API gráficos (OpenGL y Direct3D)
! Manejo de texturas:
V. Compilación
! cudaBindTexture()
! cudaBindTextureToArray()
! Interoperabilidad con los API gráficos:
! Ejemplo con OpenGL: cudaGLMapBufferObject()
! Ejemplo con Direct3D: cudaD3D9MapVertexBuffer()
! Vertex Buffer Objects y Pixel Buffer Objects, tanto de
OpenGL como de Direct3D, pueden mapearse sobre el
espacio de direcciones de CUDA, y luego se manejan
como memoria global en el código de la GPU.
! Los datos de una imagen alojados en Pixel Buffer Objects
pueden mostrarse con glDrawPixels/glTexImage2D.
83
84
82
El proceso de compilación
void serial_function(… )
{
...
}
void other_function(int ... )
{
...
}
void saxpy_serial(float ... )
{
for (int i = 0; i < n; ++i)
y[i] = a*x[i] + y[i];
}
void main( ) {
float x;
saxpy_serial(..);
...
}
Kernels
CUDA
NVCC
(Open64)
Los diferentes módulos de compilación
! El código fuente CUDA se
compila con NVCC.
Resto del
código C
! NVCC separa el código que
se ejecuta en CPU del que lo
hace en GPU.
Compilador
de la CPU
! La compilación se realiza
en dos etapas:
Identificar
los kernels
Ficheros objeto
CUDA y rees- Ficheros objeto
de la CPU
CUDA
Enlazador
cribirlos para
aprovechar
paralelismo
en GPU
Ejecutable
! Virtual: Genera código PTX
(Parallel Thread eXecution).
! Física: Genera el binario para
una GPU específica (o incluso
para una CPU multicore - ver
un poco más adelante).
CPU-GPU
C/C++ CUDA
Application
Código
fuente
NVCC
Código CPU
Virtual
PTX Code
Físico
PTX to Target
Compiler
G80
…
GPU
85
Código
objeto
86
Compilador NVCC y máquina virtual PTX
NVCC (Nvidia CUDA Compiler)
C/C++ CUDA
Application
! NVCC es un driver del compilador.
float4 me = gx[gtid];
me.x += me.y * me.z;
! EDG
! Funciona invocando todos los compiladores y herramientas
necesarias como cudacc, g++, cl, ...
! Separa código GPU y CPU.
! NVCC produce como salida:
EDG
! O bien código C para la CPU, que debe luego compilarse con el
resto de la aplicación utilizando otra herramienta.
! O bien directamente código objeto PTX.
CPU Code
! Open64
! Genera ensamblador PTX.
! Parallel Thread eXecution
(PTX)
! El ejecutable CUDA usa dos librerías dinámicas:
! Máquina virtual e ISA.
! Modelo de programación.
! Recursos y estado de ejecución.
Open64
! The CUDA runtime library (cudart)
! The CUDA core library (cuda)
PTX Code
ld.global.v4.f32
mad.f32
87
{$f1,$f3,$f5,$f7}, [$r9+0];
$f1, $f5, $f3, $f1;
88
Depuración: Modo de emulación del
dispositivo (Device Emulation Mode)
! Un ejecutable compilado en modo de emulación del
dispositivo (nvcc -deviceemu) corre íntegramente en la
CPU (host) usando el soporte en tiempo de ejecución de
CUDA sin necesidad de GPU ni driver.
! Posibilidades de modo de emulación:
VI. Depuración y optimización
! Utilizar el soporte para la depuración de que disponga la CPU
(breakpoints, inspección de registros, etcétera).
! Acceder a los datos de la GPU desde el código de la CPU.
! Llamar a cualquier función de la CPU desde el código de la GPU
(por ejemplo, printf) y viceversa.
! Detectar situaciones de interbloqueo por el uso indebido de
__syncthreads.
90
89
Carencias del modo de
emulación del dispositivo
CUDA Visual Profiler
! Los hilos del código se ejecutan de forma secuencial, por lo
que el acceso simultáneo a la misma posición de memoria
desde múltiples hilos produce (potencialmente) resultados
distintos.
! El acceso a valores mediante punteros de GPU en la CPU o
punteros de CPU en la GPU puede producir resultados
correctos en el modo de emulación, pero dará un error al ser
ejecutado.
! Los resultados de la computación en punto flotante diferirán
ligeramente debido a:
! Proporciona ciertas mediciones estratégicas para la
localización de problemas de rendimiento.
! Temporización entre GPU y CPU para las invocaciones a los kernels
y las llamadas a memcpy.
! Evolución a determinados pasos de tiempo.
! Acceso a contadores de rendimiento hardware.
! Las diferentes salidas del compilador.
! Diferentes conjuntos de instrucciones.
! El uso de precisión extendida para la generación de resultados
intermedios.
91
92
Contadores que proporciona el profiler
Interpretación de los contadores del profiler
! Representan los eventos dentro de un warp de hilos.
! Involucran a un solo multiprocesador
! Se sigue la pista de ciertos eventos definidos por el programa a través de contadores
hardware que registran determinadas señales emitidas por el chip:
!
!
!
!
gld_incoherent
gld_coherent
gst_incoherent
gst_coherent
! Los valores no se corresponderán con el número total de warps
lanzados para un kernel dado.
! Debemos lanzar un número suficiente de bloques de hilos para
asegurar que el multiprocesador involucrado tiene asignado un
porcentaje consistente del trabajo total.
Cargas/almac. en memoria global son coherentes
(“coalesced”) o incoherentes (“non-coalesced”)
! local_load
! local_store
Cargas y almacenamientos locales
! branch
! divergent_branch
Número total de bifurcaciones (“branches”) y
bifurcaciones divergentes tomadas por los hilos
! instructions – cuenta del número de instrucciones ejecutadas.
! warp_serialize – warps de hilos que se han sido secuencializados por los conflictos
de direccionamiento a la memoria compartida o a la memoria de constantes.
! cta_launched – bloques de hilos que han sido ejecutados.
! Los valores se utilizan principalmente para cuantificar la
mejora de rendimiento producida por una versión
optimizada del código.
! Se trata de reducir gld/gst_incoherent, divergent_branch
y warp_serialize.
93
94
Maximizar la ocupación
para ocultar la latencia
Restricciones en la configuración de la
ejecución
! Dónde se producen las latencias:
! Máximo nº de hilos por bloque: 512
! Nº de hilos activos limitados por los recursos.
! En el acceso a memoria global: 400-600 ciclos.
! En dependencias read-after-write entre registros.
! Nº de registros por multiprocesador: 8192 ó 16384 (serie 10 en
adelante).
! Cantidad de memoria compartida por multiprocesador: 16K ó 48K (Fermi
en adelante).
! El resultado de las instrucciones sólo puede leerse 11 ciclos más tarde.
! Se bloquean las instrucciones que presentan
dependencias en el mismo hilo, pero no las que se
encuentran en otros hilos. Para ocultar esta latencia,
debemos ejecutar tantos hilos por multiprocesador como
sea posible, eligiendo los parámetros de ejecución que
maximicen:
! Utilizar el flag–maxrregcount=N en NVCC
! N = nº máximo de registros deseado por kernel.
! En algún momento puede sobrevenir un “spilling” a LMEM (trasvase
entre el banco de registros y la memoria local que emplea el compilador
para su uso interno), con el consiguiente impacto en el rendimiento.
ocupación = (nº warps activos)/(máx. nº warps activos)
! Rastrear el archivo .cubin para descubrir el uso de LMEM
(ver la siguiente diapositiva).
Nota: El máximo nº de warps activos en las series 8, 9 y 10 de GeForce es de 24, y en
Fermi es de 48.
95
96
Heurísticos para la configuración
de la ejecución
Para conocer la utilización de los recursos
! Compilar el código del kernel con el flag -cubin para
conocer cómo se están usando los registros.
! Abrir el archivo .cubin con un editor de texto y mirar la
sección “code”.
! El número de hilos por bloque debe ser un múltiplo de 32.
! Para no desperdiciar en la ejecución de warps incompletos
! El número de bloques debe superar al de multiprocs.
! Para que todos ellos tengan al menos un bloque que ejecutar.
! Los recursos por bloque (registros y memoria compartida)
deben ser al menos la mitad del total disponible.
architecture {sm_10}
Memoria local para cada hilo
abiversion {0}
(usada por el compilador para
modname {cubin}
volcar contenidos de los registros
code {
en memoria)
name = BlackScholesGPU
lmem = 0
Memoria compartida usada
smem = 68
por cada bloque de hilos
reg = 20
bar = 0
Registros usados
bincode {
0xa0004205 0x04200780 0x40024c09 0x00200780 por cada hilo
…
! De lo contrario, resulta mejor fusionar bloques.
! El número de bloques debe superar en más del doble al
de multiprocesadores.
! Para tener siempre más de un bloque activo por multiprocesador que
garantice su ocupación en caso de que alguno se bloquee.
! El número de bloques debe superar el centenar.
! Para que el código sea escalable en futuras generaciones.
97
Heurísticos para la configuración
de la ejecución (cont.)
98
Parametrización de una aplicación
! Conflicto: Más hilos por bloque significa mejor ocultación
de latencia, pero menos registros por hilo.
! Sugerencia: Utilizar un mínimo de 64 hilos por bloque, o
incluso mejor, 192 ó 256 hilos (si aún se dispone de
registros suficientes para ello).
! Conflicto: Incrementar la ocupación no significa
necesariamente aumentar el rendimiento, pero una baja
ocupación del multiprocesador no permite ocultar latencias
en kernels limitados por el ancho de banda a memoria.
! Sugerencia: Prestar atención a la intensidad aritmética y al
paralelismo disponible.
! Todo lo que concierne al rendimiento es dependiente de la
aplicación, por lo que hay que experimentar con ella para
lograr resultados óptimos.
! Las GPUs evolucionan en muchos aspectos:
! El número de multiprocesadores. 16 en la serie 8, 30 en la serie 10,
16 de nuevo en Fermi (pero con 32 cores cada uno).
! Ancho de banda con memoria: En torno a 100 GB/s.
! Tamaño de la memoria compartida: 16 KB. para cada
multiprocesador, salvo en Fermi, que puede ampliarse a 48 KB.
! Tamaño del banco de registros: 8192 en la serie 8, 16384 en la
serie 10, 4096 en Fermi.
! Hilos: Ahora, hasta 512 por bloque. 768, 1024 ó 1536 en total
según tengamos una serie 8, una serie 10 o una arquitectura Fermi.
99
100
CUDA Occupancy Calculator
! Asesora en la selección de los parámetros de configuración
VII. Estrategias de mejora
101
102
(1) Solapamiento de
computación y comunicaciones
(1) Solapamiento de
computación y comunicaciones (cont.)
! Posibilidades:
! Ejemplo 1 (notar que 0 es el stream por defecto):
! Solapar computación en CPU con transferencias de datos a la GPU.
Es posible en cualquier dispositivo CUDA, aprovechando que el
lanzamiento de kernels desde la CPU es asíncrono. Sencillo de
implementar.
! Solapar la computación de un kernel en GPU con transferencias de
datos a la GPU. Sólo es posible a partir de CUDA Compute Cap. 1.1.
Más complejo de implementar, pues necesitamos tres cosas:
cudaMemcpyAsync(A_d, A_h, size, cudaMemcpyHostToDevice, 0);
kernel<<<grid, block>>>(A_d);
cpuFunction();
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaMemcpyAsync(dst, src, size, dir, stream1);
kernel<<<grid, block, 0, stream2>>>(…);
cudaMemcpyAsync(dst2, src2, size, stream2);
cudaStreamQuery(stream2);
103
solape
! Ejemplo 2 (las dos últimas sentencias no suponen solape adicional):
1. Alojar la memoria pinned en CPU (para que no pagine) usando
cudaMallocHost().
2. Definir streams. Un stream es una secuencia de operaciones CUDA que se
ejecutan de forma ordenada.
3. Transferir los datos CPU ! GPU usando las variantes asíncronas de copia de
datos (cudaMemcpyAsync(dst,src,size,dir,stream)), que devuelven
inmediatamente el control a la CPU.
solape
! http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls
104
(3) Optimizar los patrones
de acceso a memoria
(2) Optimizar el uso de la memoria
! Minimizar las transferencias entre CPU y GPU.
! El ancho de banda efectivo puede variar un orden de
magnitud dependiendo del patrón de acceso si sabemos
utilizar las siguientes armas:
! Recordar que este ancho de banda es muy inferior al ancho de
banda de la memoria de vídeo.
! En los casos en que se quiera aumentar el ancho de banda, utilizar
memoria “pinned” (sin abusar), que aprovecha casi todas las
prestaciones de PCI-express (3.2 GB/s. en v1.0 y 5.2 GB/s. en v2.0
normalmente, 4 y 8 GB/s. si usamos el chipset nForce 680i).
! Accesos coalesced a memoria global (menos importante en Fermi).
! Accesos a memoria compartida sin conflictos a sus 16 bancos.
! Accesos a memoria de texturas (que pasa por caché).
! Accesos a memoria de constantes que tienen una misma dirección.
! Agrupar las transferencias de datos entre CPU y GPU
! Recordar que:
! Ya que la latencia predomina en el coste sobre el ancho de banda.
! Procesar datos es más rápido que moverlos, ya que las GPUs
dedican muchos más transistores a las ALUs que a la memoria.
! Cuanto menos estrangulado se encuentre un kernel por el acceso a
memoria, mejor se comportará en las arquitecturas GPU futuras.
! Pasar algunas funciones de CPU a GPU aunque no puedan
explotar mucho paralelismo
! Si eso evita un doble trasiego de datos de GPU a CPU y regreso.
105
(5) Eliminar los conflictos en el acceso a
los 16 bancos de memoria compartida
(4) Uso de la memoria compartida
Si el tamaño se conoce
en tiempo de compilación
__global__ void kernel (...)
{
...
__shared__ float sData[256];
...
}
void main()
{
...
kernel<<<nBlocks,blocksize>>>(...);
...
}
106
! Se dispone de un espacio total de
16 Kbytes en cada multiprocesador.
! Se estructura en 16 bancos
con palabras de 32 bits
y entrelazado de orden inferior.
! El acceso simultáneo de los hilos a
memoria en cada medio warp se
produce en cualquiera de las dos
situaciones siguientes:
1. Siempre que los hilos accedan
a bancos diferentes (cualquier
permutación).
2. Si todos los hilos acceden a la
misma palabra de un banco.
Si el tamaño se desconoce
en tiempo de compilación
__global__ void kernel (...)
{
...
extern __shared__ float sData[];
...
}
void main()
{
...
smBytes = blockSize * sizeof(float);
kernel<<<nBlocks,blocksize,smBytes>>>(...);
...
}
107
Caso 1:
Caso 2:
108
(5) Eliminar los conflictos en el acceso a
los 16 bancos de memoria compartida (cont.)
! Esta optimización puede resultar bastante desafiante, aunque existe un
truco para llevarla a cabo de forma general ajustando los índices de acceso
al ID de cada grupo de trabajo:
VIII. Ejemplos
__shared float mem[N];
size_t localIdx = get_local_id[1];
for (int=0; i<N; i++)
float val = mem[i+localIdx % N];
! El precio a pagar respecto a una optimización manual es la pérdida por
computar la aritmética módulo, aunque si N es potencia de 2, se puede
reemplazar por un operador AND bit-wise mucho más eficiente:
assert (N is a power of 2);
__shared float mem[N];
size_t localIdx = get_local_id[1];
for (int=0; i<N; i++)
float val = mem[i+localIdx & (N-1)];
109
110
Pasos a realizar para
la construcción del código
1. Identificar las partes potencialmente paralelas.
2. Delimitar los datos necesarios.
3. Mover los datos a la GPU.
4. Llamar al kernel de computación.
5. Establecer la sincronización apropiada entre GPU y CPU.
6. Transferir los resultados de la GPU a la CPU.
7. Integrar los resultados de la GPU en variables de la CPU.
VIII.1. Suma de dos vectores
111
112
Código necesario para
el kernel y su invocación
// Suma de los vectores C = A+B
// Cada hilo calcula un componente del vector
__global__ void vecAdd(float* A, float* B, float* C)
{
!
Int tid = threadIdx.x+ blockDim.x* blockIdx.x;
!
C[tid] = A[tid] + B[tid];
}
Código en CPU para el manejo de memoria
y la recolección de resultados desde la GPU
Código GPU
int main()
{ // Lanzamos N/256 bloques de 256 hilos cada uno
!
vecAdd<<< N/256, 256>>>(d_A, d_B, d_C);
Código CPU
}
113
unsigned int numBytes = N * sizeof(float);
// aloja memoria en la CPU
float* h_A = (float*) malloc(numBytes);
float* h_B = (float*) malloc(numBytes);
... inicializa h_A y h_B ...
// aloja memoria en la GPU
float* d_A = 0; cudaMalloc((void**)&d_A, numbytes);
float* d_B = 0; cudaMalloc((void**)&d_B, numbytes);
float* d_C = 0; cudaMalloc((void**)&d_C, numbytes);
// copia los datos de entrada de la CPU en la GPU
cudaMemcpy(d_A, h_A, numBytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, numBytes, cudaMemcpyHostToDevice);
... LA LLAMADA AL KERNEL vecAdd ANTERIOR SE REALIZA AQUI ...
// copia los resultados de la GPU en la CPU
float* h_C = (float*) malloc(numBytes);
cudaMemcpy(h_C, d_C, numBytes, cudaMemcpyHostToDevice);
// libera la memoria de vídeo
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
64
114
Resumen de actuaciones a realizar para
mejorar el rendimiento del código CUDA
CUDA permite lograr excelentes resultados en la
paralelización de código si respetamos unas pocas
premisas de rendimiento:
VIII. Producto de matrices
1. Estructurar el código y elegir una configuración de hilos que
maximice el paralelismo de datos en GPU y minimice la
transferencia de datos CPU ! GPU.
2. Respetar la coherencia al acceder a memoria global (coalescing).
3. Maximizar el uso de la memoria compartida (tiling).
4. Evitar accesos a memoria compartida con elevado número de
conflictos al acceder a sus bancos.
5. Minimizar el uso de warps divergentes.
Vamos a aplicar 1, 2 y 3 sobre el producto de matrices.
115
116
Optimización 1: Estructurar el código
para maximizar el paralelismo
Optimización 2: Coalescing para
maximizar el acceso a memoria global
! Declarar muchos threads permite ocultar latencias con memoria global.
! Máxima granularidad del código: Utilizar un único grid con tantos hilos
como elementos haya en la matriz C, de forma que cada hilo computa un
único elemento de la matriz resultado C.
! Entre el grid y los hilos se sitúa el bloque para acomodar los hilos en
grupos de a lo sumo 512 threads (máx. paralelismo en CUDA). Y como 512
no es un cuadrado perfecto, 256 resulta el candidato ideal (256 = 16x16)
para equilibrar el paralelismo en las dos dimensiones de la matriz.
Grid
WidthA
WidthB
Bloque
····
······ !!!!!!!
····
····
······
····
!!!!!!!!!!!!!!!
····
!!!!!!!
· · · · · · Th(x,y)
····
····
······
····
! Las instrs. pueden leer/escribir palabras de 32, 64 o 128 bits.
! Medio warp puede acceder de forma unificada (coalesced) para
leer 32, 64 o 128 bytes a segmentos de memoria global.
! Segmento: Intervalo continuo de memoria de 32 , 64 o 128 bytes,
alineado a estos tamaños.
K mod 32
WidthB
32 bytes
K mod 128
K mod 64
64 bytes
HeightA
C(x, y)
C
=
HeightA
X
A
128 bytes
B
dim2 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
dim2 dimGrid(WidthB/BLOCKSIZE, HeightA/BLOCKSIZE);
Mem. global
Mem. global
Mem. global
117
Coalescing (3)
Coalescing (2)
! Condiciones para el coalescing de hilos en CUDA 1.0 y 1.1
! Condiciones de coalescing para CUDA 1.2 o
superior:
! Acceso a palabras de 32 bits (un acceso de 64 bytes), de 64 bits (dos accesos de
128 bytes) o de 128 bits (cuatro accesos de 128 bytes)
! Las 16 palabras deben estar ubicadas en un segmento de memoria de igual tamaño
al de la transacción (o del doble en el caso de 128 bits)
! Los hilos deben acceder en secuencia (el k-ésimo hilo a la k-ésima palabra)
! Los hilos acceden a datos ubicados en un segmento
de memoria de:
! 32 bytes (datos de 8 bits).
! 64 bytes (datos de 16 bits).
! 128 bytes (datos de 64 bits).
! Si las condiciones no se cumplen, cada hilo requerirá un nuevo acceso.
Coalesced
118
No-coalesced
! El patrón de acceso puede ser cualquiera, e incluso
varios hilos pueden acceder al mismo dato.
! El número simultáneo de accesos a memoria, n,
será aquel que pueda acomodarse en medio warp
conformando n segmentos.
119
120
Optimización 2: Ejemplo de coalescing sobre
el producto de matrices
Optimización 3: Tiling aplicado
al producto de matrices
! C = A ! B. Matrices de N x N datos.
! Sin tiling:
Coalescing en acceso a la matriz B:
B
! Requiere un gran
ancho de banda.
! Tiene poca
intensidad
aritmética.
Memoria global
! Cada hilo necesita 10 regs., por lo que podemos declarar
768 hilos estructurados en 3 bloques de 256 (16x16).
Rendimiento: 10.58 GFLOPS
Mirando el PTX del código, el cuello de botella se sitúa en
el acceso a memoria global.
A
N
! Cada hilo computa un elemento de C.
! A y B se cargan N veces de memoria de vídeo.
C
N
A
Ctemp=0;
for (i=0; i<widthA; i++){
Ctemp += A[indexA] * B[indexB];
indexA++;
indexB += widthB;
}
C[indexC] = Ctemp ;
B
N
N
121
Optimización 3: Tiling aplicado
al producto de matrices (2)
Optimización 3: Implementación del tiling
B
Tenemos que particionar los datos en subconjuntos (tiles) que
quepan en los 16 Kbytes de memoria compartida…
… y manipular cada subconjunto desde un bloque de hilos:
M
! Cada bloque de hilos se encarga de una
submatriz Csub de MxM elementos de C.
N
M
! Con tiling:
! A y B se cargan sólo (N/M) veces desde
memoria de vídeo.
! Cargar el subconjunto desde memoria global a memoria compartida.
! __syncthreads()
! Realizar la computación del subconjunto en memoria compartida.
M
A
C
! Cada hilo puede iterar eficientemente sobre cualquier elemento de datos.
! __syncthreads() (en caso de necesidad por dependencias)
! Copiar los resultados desde memoria compartida a memoria global.
M
M
N
M
N
Csub
M
! Relaja el ancho
de banda.
! Aumenta la
intensidad
aritmética.
122
M
N
123
124
Optimizaciones adicionales:
Loop unrolling
Optimización 3: Implementación del tiling
GRID
!!!!!!!!!!
SM0
As
!!!!
!!!!
!!!!
Bs
!!!!
!!!!
!!!!
8192 registros
SP0
SM15
Shared Memory
8 KB constant
!!!!!!!!!!!!!!
SFU0
SP7
SFU1
! Cada hilo carga un elemento
del tile de A y B
¡El compilador también puede hacer cosas por nosotros!
Ctemp=0;
for (!!!){
__shared__ float As[16][16];
__shared__ float Bs[16][16];
// Load tile (16x16)
As[ty][tx] = A[indexA];
Bs[ty][tx] = B[indexB];
indexA += 16;
indexB += 16 * widthB;
__syncthreads();
// Compute results from tile
for (i=0; i<16; i++)
Ctemp+=As[ty][i]*Bs[i][tx];
__syncthreads();
}
C[indexC] = Ctemp ;
Ctemp=0;
for (!!!){
__shared__ float As[16][16];
__shared__ float Bs[16][16];
// Load tile (16x16)
As[ty][tx] = A[indexA];
Bs[ty][tx] = B[indexB];
indexA += 16;
indexB += 16 * widthB;
__syncthreads();
// Compute results from tile
for (i=0; i<16; i++)
Ctemp+=As[ty][i]*Bs[i][tx];
__syncthreads();
}
C[indexC] = Ctemp ;
Ctemp=0;
for (!!!){
__shared__ float As[16][16];
__shared__ float Bs[16][16];
// Load tile (16x16)
As[ty][tx] = A[indexA];
Bs[ty][tx] = B[indexB];
indexA += 16;
indexB += 16 * widthB;
__syncthreads();
// Compute results from tile
Ctemp+=As[ty][0]*Bs[0][tx];
!!!!
Ctemp+=As[ty][15]*Bs[15][tx];
__syncthreads();
}
C[indexC] = Ctemp ;
!!!!
!!!!
!!!!
!!!!
!!!!
!!!!
A
B
Memoria global
125
126
Rendimiento de tiling & unrolling
VIII. 3. Incremento de un escalar
a un vector de float3
100
75
GFLOPS
50
Sólo tiling
Tiling & Unrolling
25
0
4x4
8x8
12x12
16x16
Tamaño del tile
127
128
Dónde está el problema: Coalescing
Código de partida
! float3 ocupa 12 bytes.
! Cada hilo concluye ejecutando tres lecturas.
__global__ void accessFloat3(float3 *d_in, float3 d_out)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
float3 a = d_in[index];
! sizeof(float3) ! 4, 8, o 16.
! Cada half-warp lee tres regiones no-contiguas de 64 bytes.
a.x += 2;
a.y += 2;
a.z += 2;
d_out[index] = a;
t0
t1
t2
float3
float3
float3
t3
! Debemos cambiar el patrón de acceso para que los hilos
accedan a memoria global de forma consecutiva.
}
129
Ideas para llevar a cabo
la implementación
Cómo resolver el problema
MEMORIA GLOBAL
…
Paso 2
Paso 1
…
t0 t1 t2
t255
…
…
MEMORIA COMPARTIDA
130
t0 t1 t2
…
…
…
MEMORIA COMPARTIDA
! Usar la memoria compartida para permitir coalescing
! Se necesitan sizeof(float3)*(hilos/bloque) bytes de memoria compartida.
! Cada hilo lee tres escalares de tipo float:
! Desplazamiento: 0, (hilos/bloque), 2*(hilos/bloque).
! Estos datos serán presumiblemente procesados por otros hilos, por lo que se
requiere sincronizar.
! Procesamiento
! Cada hilo recupera su float3 del vector en memoria compartida
! Para ello, hay que convertir el puntero a memoria en (float3*) y utilizar threadID
como índice.
! El resto del código no requiere modificación alguna.
De forma similar, el paso 3 comienza
desplazado 512 posiciones.
131
132
Codificación en CUDA
Leer operandos
fuente a través
de memoria
compartida
Rendimiento obtenido
__global__ void accessInt3Shared(float *g_in, float *g_out)
{
int index = blockIdx.x * blockDim.x + threadIdx.x;
__shared__ float s_data[256*3];
s_data[threadIdx.x]
= g_in[index];
s_data[threadIdx.x+256] = g_in[index+256];
s_data[threadIdx.x+512] = g_in[index+512];
__syncthreads();
float3 a = ((float3*)s_data)[threadIdx.x];
Este código
permanece
inalterado
a.x += 2;
a.y += 2;
a.z += 2;
Escribir el
resultado a
través de
memoria
compartida
((float3*)s_data)[threadIdx.x] = a;
__syncthreads();
g_out[index]
= s_data[threadIdx.x];
g_out[index+256] = s_data[threadIdx.x+256];
g_out[index+512] = s_data[threadIdx.x+512];
}
! Condiciones experimentales:
! Creamos un kernel que lee un array de floats, incrementa sus
posiciones y escribe el resultado.
! Utilizamos 3 millones de floats (12 Mbytes)
! Promediamos los tiempos sobre 10000 ejecuciones.
! Para float: 12K bloques de 256 hilos.
! 356"s – coalesced.
! 3,494"s – acceso permutado/desalineado por cada hilo.
! Para float3: 4K bloques de 256 hilos:
! 3,302"s – NO coalesced.
! 359"s – Coalesced a través de memoria compartida.
133
134
Bibliografía
Herramientas software
! GPU Gems, Gems 2 y Gems 3.
! CUDA Programming Guide. Las bases de CUDA.
! CUDA Best Practices Guide. Para optimizar código.
! CUDA Zone (http://www.nvidia.com/cuda).
! CUDA Toolkit, ya en su versión 3.0 para Fermi.
! Compilador nvcc, librerías y documentación.
! Implementaciones disponibles para Linux, MacOS y Windows 7.
! CUDA SDK:
! Scripts de compilación y algunas utilidades.
! Códigos de ejemplo y whitepapers.
! Los códigos que se han desarrollado en CUDA junto a los factores
de aceleración logrados.
! Los artículos de investigación que describen las aplicaciones y su
implementación.
! Tutoriales, forums, curso de programación paralela de la UIUC, ...
! GPGPU: El movimiento originario de la programación de
propósito general sobre la GPU que desde 2003 aglutina
las novedades más interesantes. http:///www.gpgpu.org
! CUDA Occupancy Calculator.
! CUDA Profiler: Analiza tiempos de ejecución, accesos
coalesced a memoria, warp divergentes, conflictos en
el acceso a memoria compartida, ...
! Nvidia Nexus: Entorno de desarrollo similar a Visual Studio.
135
136