PDF de programación - Programando la GPU con CUDA

Río Cuarto (Argentina), 17 y 18 de Febrero, 2014

RIO 2014

Agradecimientos

Al personal de Nvidia, por compartir conmigo ideas,

material, diagramas, presentaciones, ... Alfabéticamente:

Bill Dally [2010-2011: Consumo energético, Echelon y diseños futuros].
Simon Green [2007-2009: Los pilares de CUDA].
Sumit Gupta [2008-2009: El hardware de Tesla].
Mark Harris [2008, 2012: CUDA, OpenACC, Lenguajes de Programación, Librerías].
Wen-Mei Hwu [2009: Programación y trucos para mejorar el rendimiento].
Stephen Jones [2012: Kepler].
David B. Kirk [2008-2009: Hardware de Nvidia].
Timothy Lanfear [2012: Kepler y OpenACC].
David Luebke [2007-2008: Hardware de Nvidia].
Lars Nyland [2012: Kepler].
Edmondo Orlotti [2012: CUDA 5.0, OpenACC].
Cyril Zeller [2008-2009: Fundamentos de CUDA].

... por mencionar sólo unos pocos del equipo que contribuyó a esta

presentación.

Prerrequisitos para este tutorial

Se requiere estar familiarizado con el lenguaje C.
No es necesaria experiencia en programación paralela

(aunque si se tiene, ayuda bastante).

No se necesitan conocimientos sobre la arquitectura de la

GPU: Empezaremos describiendo sus pilares básicos.

No es necesario tener experiencia en programación

gráfica. Eso era antes cuando GPGPU se basaba en los
shaders y Cg. Con CUDA, no hace falta desenvolverse con
vértices, píxeles o texturas porque es transparente a todos
estos recursos.

2

4

Manuel Ujaldon
Nvidia CUDA Fellow
Profesor del Dpto. de Arquitectura de Computadores. Universidad de Málaga (España).
Conjoint Senior Lecturer. University of Newcastle (Australia)

Contenidos del tutorial

1. Introducción. [18 diapositivas]
2. Arquitectura. [15]

1. El modelo hardware de CUDA. [3]
2. La primera generación: Tesla (2006-2009). [3]
3. La segunda generación: Fermi (2010-2011). [4]
4. La tercera generación: Kepler (2012-2014). [5]

3. Programación. [15]
4. Sintaxis. [16]

1. Elementos básicos. [10]
2. Un par de ejemplos preliminares. [6]

5. Compilación. [9]
6. Ejemplos: VectorAdd, Stencil, MxM. [25]
7. Bibliografía y herramientas. [6]

3

Programando la GPUcon CUDA Bienvenido al mundo de las GPUs

Los personajes de esta historia:
La foto de familia de CUDA

En apenas 5 años, la programación CUDA ha
crecido a un ritmo vertiginoso. Así, en 2013:

Se han publicado más de 40.000 artículos científicos.
CUDA se enseña en más de 650 universidades.
Hay más de 450 millones de GPUs programables con CUDA.
Más de 150.000 desarrolladores/programadores en activo.
Más de 1.700.000 descargas del compilador y SDK.

7

6

8

I. Introducción Las 3 cualidades que han hecho
de la GPU un procesador único

¿Qué es CUDA?
“Compute Unified Device Architecture”

Control simplificado.

El control de un hilo se amortiza en otros 31 (warp size = 32).

Escalabilidad.

Aprovechándose del gran volumen de datos que manejan las

aplicaciones, se define un modelo de paralelización sostenible.
Productividad.

Se habilitan multitud de mecanismos para que cuando un hilo pase

a realizar operaciones que no permitan su ejecución veloz, otro
oculte su latencia tomando el procesador de forma inmediata.

Palabras clave esenciales para CUDA:

Warp, SIMD, ocultación de latencia, conmutación de contexto

gratis.

Lo esencial de CUDA

En general, es lenguaje C con mínimas extensiones:

El programador escribe el programa para un solo hilo (thread), y el

código se instancia de forma automática sobre miles de hilos.
CUDA define:

Un modelo de arquitectura:

Con multitud de unidades de proceso (cores), agrupadas en multiprocesadores que

comparten una misma unidad de control (ejecución SIMD).

Un modelo de programación:

Basado en el paralelismo masivo de datos y en el paralelismo de grano fino.
Escalable: El código se ejecuta sobre cualquier número de cores sin recompilar.

Un modelo de gestión de la memoria:

Más explícita al programador, con control explícito de la memoria caché.

Objetivos:

Construir código escalable a cientos de cores, declarando miles de hilos.
Permitir computación heterogénea en CPU y GPU.

9

11

Una plataforma diseñada conjuntamente a nivel software y

hardware para aprovechar la potencia de una GPU en
aplicaciones de propósito general a tres niveles:

Software: Permite programar la GPU en C con mínimas

pero potentes extensiones SIMD para lograr una ejecución
eficiente y escalable.

Firmware: Ofrece un driver para la programación GPGPU

que es compatible con el que utiliza para renderizar. Sencillos
APIs manejan los dispositivos, la memoria, etc.

Hardware: Habilita el paralelismo de la GPU para

programación de propósito general a través de un número de
multiprocesadores dotados de un conjunto de núcleos
computacionales arropados por una jerarquía de memoria.

Computación heterogénea (1/4)

Terminología:

Host (el anfitrión): La CPU y la memoria de la placa base [DDR3].
Device (el dispositivo): La tarjeta gráfica [GPU + memoria de vídeo]:

GPU: Nvidia GeForce/Tesla.
Memoria de vídeo: GDDR5 en 2013.

Host

Device

10

12

Computación heterogénea (2/4)

Computación heterogénea (3/4)

CUDA ejecuta un programa sobre un dispositivo (la GPU), que actúa como

coprocesador de un anfitrión o host (la CPU).

CUDA puede verse como una librería de funciones que contienen 3 tipos de

componentes:

Host: Control y acceso a los dispositivos.
Dispositivos: Funciones específicas para ellos.
Todos: Tipos de datos vectoriales y un conjunto de rutinas soportadas por ambas partes.

CPU (host)

Cores

Caches

50 GB/s.

GPU
(device)

3 canales (192 bits = 24 bytes)

@ 1.333 GHz 32 GB/s.

384 bits @ 3 GHz 144 GB/s.

Memoria principal

(DDR3)

PCI-e 3.0: 8 GB/s.

Memoria de

vídeo

(GDDR5)

El código reescrito en CUDA puede ser inferior al 5%, pero

13

consumir más del 50% del tiempo si no migra a la GPU.

14

Computación heterogénea (4/4)

Un sencillo flujo de procesamiento (1/3)

#include <iostream>
#include <algorithm>

using namespace std;

#define N 1024
#define RADIUS 3
#define BLOCK_SIZE 16

__global__ void stencil_1d(int *in, int *out) {



__shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
int gindex = threadIdx.x + blockIdx.x * blockDim.x;
int lindex = threadIdx.x + RADIUS;



}

// Read input elements into shared memory
temp[lindex] = in[gindex];
if (threadIdx.x < RADIUS) {


}

temp[lindex - RADIUS] = in[gindex - RADIUS];
temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];

// Synchronize (ensure all the data is available)
__syncthreads();

// Apply the stencil
int result = 0;
for (int offset = -RADIUS ; offset <= RADIUS ; offset++)


result += temp[lindex + offset];

// Store the result
out[gindex] = result;

void fill_ints(int *x, int n) {

}

fill_n(x, n, 1);

int main(void) {



int *in, *out; // host copies of a, b, c
int *d_in, *d_out; // device copies of a, b, c
int size = (N + 2*RADIUS) * sizeof(int);



}

// Alloc space for host copies and setup values
in = (int *)malloc(size); fill_ints(in, N + 2*RADIUS);
out = (int *)malloc(size); fill_ints(out, N + 2*RADIUS);

// Alloc space for device copies
cudaMalloc((void **)&d_in, size);
cudaMalloc((void **)&d_out, size);

// Copy to device
cudaMemcpy(d_in, in, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_out, out, size, cudaMemcpyHostToDevice);

// Launch stencil_1d() kernel on GPU
stencil_1d<<<N/BLOCK_SIZE,BLOCK_SIZE>>>(d_in + RADIUS, d_out + RADIUS);

// Copy result back to host
cudaMemcpy(out, d_out, size, cudaMemcpyDeviceToHost);

// Cleanup
free(in); free(out);
cudaFree(d_in); cudaFree(d_out);
return 0;

CODIGO DEL DISPOSITIVO:
Función paralela

CODIGO DEL HOST:

- Código serie.
- Código paralelo.
- Código serie.

Bus PCI

1.Copiar los datos de entrada de la
memoria de la CPU a la memoria
de la GPU.

15

16

Un sencillo flujo de procesamiento (2/3)

Un sencillo flujo de procesamiento (3/3)

PCI Bus

PCI Bus

1.Copiar los datos de entrada de la
memoria de la CPU a la memoria
de la GPU.

2.Cargar el programa en la GPU y

ejecutarlo, ubicando datos en
caché para mejorar el
rendimiento.

El clásico ejemplo

int main(void) {

printf("¡Hola mundo!\n");
return 0;

}

Salida:

$ nvcc hello.cu
$ a.out
¡Hola mundo!
$

Es código C estándar que se ejecuta en el host.
El compilador nvcc de Nvidia puede utilizarse para

compilar programas que no contengan código para la GPU.

1.Copiar los datos de entrada de la
memoria de la CPU a la memoria
de la GPU.

2.Cargar el programa en la GPU y

ejecutarlo, ubicando datos en
caché para mejorar el
rendimiento.

3.Transferir los resultados de la

memoria de la GPU a la memoria
de la CPU.

17

19

18

¡Hola mundo! con código para la GPU (1/2)

__global__ void mikernel(void)

{
}

int main(void)
{

mikernel<<<1,1>>>();

printf("¡Hola mundo!\n");
return 0;

}

Dos nuevos elementos

sintácticos:

La palabra clave de CUDA __global__
indica una función que se ejecuta en la
GPU y se lanza desde la CPU. Por
ejemplo, mikernel<<<1,1>>>.
Eso es todo lo que se requiere

para ejecutar una función en GPU.

nvcc separa el código fuente para la CPU y la GPU.
Las funciones que corresponden a la GPU (como mikernel())
son procesadas por el compilador de Nvidia.
Las funciones de la CPU (como main()) son procesadas por
su compilador (como gcc para Unix o cl.exe para Windows).

20

¡Hola mundo! con código para la GPU (2/2)

Si tenemos una arquitectura CUDA, podemos
programarla de muy diversas formas...

Computación en GPU

C++

C

OpenCLtm

Direct

Compute

Fortran

Java y
Python

NVIDIA GPU con la arquitectura
de computación paralela CUDA

... aunque este tutorial se focaliza sobre CUDA C.

22

__global__ void mikernel(void)

{
}

int main(void) {

mikernel<<<1,1>>>();

printf("¡Hola mundo!\n");
return 0;

}

Salida:

$ nvcc hello.cu
$ a.out
¡Hola mundo!
$

mikernel() no hace nada esta vez.
Los símbolos "<<<" y ">>>" delimitan la llamada desde el código
de la CPU al código de l