PDF de programación - Introducción a la Computación Heterogénea - Comunicación Directa

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Introducción a la Computación Heterogénea

Comunicación Directa

Carlos Bederián, Nicolás Wolovick

FaMAF, Universidad Nacional de Córdoba, Argentina

22 de Febrero de 2013

[email protected]

Revisión 3890, 2013-02-22

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Velocidades

Todo depende por donde pasan los datos.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Copias host↔device

Para transferir por el PCIe se necesita memoria que no swapee.

Pinned Memory

Para pedir/devolver page-locked memory en CUDA:

cudaError_t cudaMallocHost(void **ptr, size_t size)

Es una versión de cudaHostAlloc sin el tercer parámetro.

cudaError_t cudaFreeHost (void* ptr)

Usan los syscalls:

#include <sys/mman.h>

int mlock(void *addr, size_t length);
int munlock(void *addr, size_t length);

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Pinned Memory – performance

Evita una copia, debería ser 2x más rápido.

“Performance of GPUDirect with MPI”, NVIDIA, 2012.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Pinned Memory – limitaciones

Es un recurso escaso: RLIMIT_MEMLOCK.

$ ulimit -a | grep locked
max locked memory

(kbytes, -l) 64

El runtime de CUDA se las arregla para utilizar este poquitito.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Copy engines

Las placas de línea Tesla permiten . . .

rio66@mini:~$ LD_LIBRARY_PATH=/opt/cuda-4.2.9/lib64/ /opt/cudasdk/C/bin/linux/release/deviceQuery
[deviceQuery] starting...

/opt/cudasdk/C/bin/linux/release/deviceQuery Starting...

CUDA Device Query (Runtime API) version (CUDART static linking)

Found 2 CUDA Capable device(s)

Device 0: "Tesla C2070"

CUDA Driver Version / Runtime Version
CUDA Capability Major/Minor version number:
Total amount of global memory:
(14) Multiprocessors x ( 32) CUDA Cores/MP:
GPU Clock rate:
Memory Clock rate:
...
Concurrent copy and execution:
...
Support host page-locked memory mapping:
Concurrent kernel execution:
...
Device supports Unified Addressing (UVA):
...

5.0 / 4.2
2.0
6143 MBytes (6441598976 bytes)
448 CUDA Cores
1147 MHz (1.15 GHz)
1494 Mhz

Yes with 2 copy engine(s)

Yes
Yes

Yes

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Superposición

• 1 kernel engine.
• 2 copias, una PCIe up engine y otra down engine.

¿Cómo puedo usar concurrentemente estas unidades?

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Superposición

• 1 kernel engine.
• 2 copias, una PCIe up engine y otra down engine.

¿Cómo puedo usar concurrentemente estas unidades?

Streams

• Cola de tareas.
• Orden secuencial intra stream.
• Absoluto no-determinismo extra streams (concurrencia!).

Cuidado, necesita que la memoria host esté pinned.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Streams

Default Stream
cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a);
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

El truco fácil para hacer superposición.

Overlap CPU-GPU
cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a);
do_something_CPU(a);
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

Cuidado, el default stream tiene propiedades especiales.

• Es mútuamente exclusiva con las otras!

“How to Overlap Data Transfers in CUDA C/C++”, ∀, 2012.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Streams – ciclo de vida

Creación
cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1)

Uso
result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1)
increment<<<1,N,0,stream1>>>(d_a)

Sincronización
result = cudaStreamSynchronize(stream0)

Destrucción
result = cudaStreamDestroy(stream1)

El tercer parámetro de <<< >>>, es la cantidad de memoria compartida dinámica.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Paralelismo al final de la cola

En Fermi:

• Los streams no son completamente paralelos.
• Serializados por una única cola de trabajos.
• Pensar esto al usar las unidades: D2H, H2D, KER.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Streams – ejemplo

cudaStream_t stream0, stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
float *d_A0, *d_B0, *d_C0; // device memory para stream 0
float *d_A1, *d_B1, *d_C1; // device memory para stream 1
// cudaMalloc ...

for (int i = 0; i < n; i += SegSize*2) {

cudaMemcpyAsync(d_A0, h_A+i, SegSize*sizeof(float),.., stream0);
cudaMemcpyAsync(d_B0, h_B+i, SegSize*sizeof(float),.., stream0);
vecAdd<<<SegSize/256, 256, 0, stream0>>>(d_A0, d_B0, d_C0);
cudaMemcpyAsync(d_C0, h_C+i, SegSize*sizeof(float),.., stream0);

cudaMemcpyAsync(d_A1, h_A+i+SegSize, SegSize*sizeof(float),.., stream1);
cudaMemcpyAsync(d_B1, h_B+i+SegSize, SegSize*sizeof(float),.., stream1);
vecAdd<<<SegSize/256, 256, 0, stream1>>>(d_A1, d_B1, d_C1);
cudaMemcpyAsync(d_C1, h_C+i+SegSize, SegSize*sizeof(float),.., stream1);

}

h2d0, h2d0, ker0, h2d0, h2d1 , h2d1, ker1, h2d1

Totalmente secuencial!

“Heterogeneous Parallel Programming”, Wen-mei W. Hwu, 2012.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Streams – ejemplo – 2

Eliminar un poco de dependencias.
for (int i = 0; i < n; i += SegSize*2) {

cudaMemcpyAsync(d_A0, h_A+i, SegSize*sizeof(float),.., stream0);
cudaMemcpyAsync(d_B0, h_B+i, SegSize*sizeof(float),.., stream0);
cudaMemcpyAsync(d_A1, h_A+i+SegSize, SegSize*sizeof(float),.., stream1);
cudaMemcpyAsync(d_B1, h_B+i+SegSize, SegSize*sizeof(float),.., stream1);
vecAdd<<<SegSize/256, 256, 0, stream0>>>(d_A0, d_B0, d_C0);
vecAdd<<<SegSize/256, 256, 0, stream1>>>(d_A1, d_B1, d_C1);
cudaMemcpyAsync(d_C0, h_C+i, SegSize*sizeof(float),.., stream0);
cudaMemcpyAsync(d_C1, h_C+i+SegSize, SegSize*sizeof(float),.., stream1);

}

Hay un poco de paralelismo.

h2d0, h2d0, h2d1, h2d1,

h2d0, h2d1

ker0, ker1,

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

HyperQ

No hay restricciones en Kepler (full): K20, K20X, GTX Titan.
HyperQ da 32 colas físicas que trabajan en paralelo.

Excelente para GPU-paralelizar (rápidamente) un programa MPI.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Work queues

• Fermi puede tener activos hasta 16 grids,

pero síncronas en inicio y fin! (especulación nuestra)

• El scheduler de Kepler (GK110) es mucho más complejo.

Mayor grado de utilización de la pastilla.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

UVA

Tenemos Unified Virtual Addressing–UVA (CUDA≥4.0)

• Punteros adornados para identificar CPU, GPU0, GPU1, . . .
• Elimina el paso por la CPU: direct access, direct transfer .

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Habilitando el acceso directo

cudaDeviceEnablePeerAccess(1,0);

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Para copiar entre placas directamente.

Direct transfer

// http://kylespafford.com/2012/02/29/gpu-direct.html
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1,flags); //flags=0
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0,flags); //flags=0

// Pedir memoria
float *gpu0data, *gpu1data;
cudaSetDevice(0);
cudaMalloc(&gpu0data, nbytes);
cudaSetDevice(1);
cudaMalloc(&gpu1data, nbytes);

// Hace la copia p2p
cudaMemcpy(gpu0data, gpu1data, cudaMemcpyDefault);

// Quitar el acceso
cudaDeviceDisablePeerAccess(0);
cudaDeviceDisablePeerAccess(1);

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

El kernel accede a la memoria de otra placa!

Direct access

__global__ void kernel(float *src, float *dst) {

const int i = blockIdx.x * blockDim.x + threadIdx.x;
dst[i] = src[i];

}

void main() {
...

cudaSetDevice(0);
kernel<<<b,t>>> (gpu0data, gpu1data);
cudaSetDevice(0);
kernel<<<b,t>>> (gpu1data, gpu0data);
cudaSetDevice(1);
kernel<<<b,t>>> (gpu0data, gpu1data);
cudaSetDevice(1);
kernel<<<b,t>>> (gpu1data, gpu0data);

Podemos hacer algo similar a OpenMP para GPUs.
(de hecho pueden interoperar bien)

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Topología y localidad

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Topología y localidad – 2

• Programas y job-scheduler, NUMA-aware.
• Se puede usar hwloc para mirar la topología.
• lstopo.

Motivación

Pinned Memory

Superposición de Comunicación y Computación

UVA

GPUDirect2

Fin

Motivación

Pinned Memory

Superposici