CUDA : [cours]    [mise en pratique]

2. Transfert Mémoire

L'objectif du TD est de comparer les temps de transfert entre mémoire de l'host et mémoire du device.

2.1. Travail à réaliser

2.1.1. Programme

Ecrire un programme CUDA qui prend en paramètre une taille de tableau en octets puis qui :

• crée un tableau d'octets sur l'host de la taille demandée
• crée un tableau équivalent sur le device

  cudaMalloc(void **devPtr, size_t size);

• transfert le tableau de l'host vers le device

  cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);

• transfert le tableau du device vers l'host

2.1.2. Présentation des résultats

On mesurera les temps de transfert en utilisant nvprof :

nvprof --unified-memory-profiling off program arguments

La sortie de nvprof se fait sur stderr, il faut récupérer les temps figurés en rouge et les convertir en secondes. Dans le cas présent il s'agit du temps moyen sur 10 itérations (Calls).

==11839== NVPROF is profiling process 11839, command: ./bin/memory_transfer.exe 1073741824
==11839== Profiling application: ./bin/memory_transfer.exe 1073741824
==11839== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 50.36%  1.06122s        10  106.12ms  105.51ms  107.42ms  [CUDA memcpy DtoH]
 49.64%  1.04620s        10  104.62ms  103.81ms  107.23ms  [CUDA memcpy HtoD]

==11839== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 95.52%  2.10938s        20  105.47ms  103.82ms  107.60ms  cudaMemcpy
  4.44%  98.097ms         1  98.097ms  98.097ms  98.097ms  cudaMalloc
  0.02%  466.87us        90  5.1870us     265ns  182.11us  cuDeviceGetAttribute
  0.01%  234.44us         1  234.44us  234.44us  234.44us  cuDeviceTotalMem
  0.00%  44.040us         1  44.040us  44.040us  44.040us  cuDeviceGetName
  0.00%  2.3910us         2  1.1950us     367ns  2.0240us  cuDeviceGetCount
  0.00%     847ns         2     423ns     316ns     531ns  cuDeviceGet

On générera un graphique en utilisant gnuplot comme ci-dessous.

2.2. Quelques résultats

Voici à présent plusieurs résultats comparatifs entre CPU et carte graphique. les colonnes représentent :

• taille : le nombre d'octets à transférer
• H2D : le temps moyen de transfert (sur 10 itérations) en micro secondes de l'host vers le device
• H2D BW: la bande passante (sur 10 itérations) en Mo/s de l'host vers le device
• D2H : le temps moyen de transfert (sur 10 itérations) en micro secondes du device vers l'host
• D2H BW: la bande passante (sur 10 itérations) en Mo/s du device vers l'host

On notera qu'il s'agit de la mémoire dite paginable (pageable) et non la mémoire pinned (non paginable).

GTX 770

GTX 970

GTX 1070

Quadro K610m

2.3. Comparaison

On pourra comparer les résultats obtenus avec l'utiliaire 1_Utilities/bandWidth des examples (samples) de NVidia.

Par example avec une GTX 970, pour transférer 16_777_216 octets, on obtient :

./bandwidthTest --memory=pageable --mode=range --start=16777216  --end=16777216  --increment=1 
[CUDA Bandwidth Test] - Starting...
Running on...

 Device 0: GeForce GTX 970
 Range Mode

 Host to Device Bandwidth, 1 Device(s)
 PAGEABLE Memory Transfers
   Transfer Size (Bytes)	Bandwidth(MB/s)
   16777216			10057.4

 Device to Host Bandwidth, 1 Device(s)
 PAGEABLE Memory Transfers
   Transfer Size (Bytes)	Bandwidth(MB/s)
   16777216			9485.5

 Device to Device Bandwidth, 1 Device(s)
 PAGEABLE Memory Transfers
   Transfer Size (Bytes)	Bandwidth(MB/s)
   16777216			138641.0

Result = PASS

Soit des résultats comparables :