Tutoriel CUDA

Introduction

Dans ce tutoriel, nous présentons le composant de calcul hautement parallèle qu'est le processeur graphique (ou GPU pour graphic processing unit) et nous abordons le modèle CUDA avec quelques-unes de ses librairies numériques utilisées en calcul de haute performance.

Prérequis

Ce tutoriel démontre comment utiliser CUDA pour accélérer des programmes en C ou en C++. Une bonne connaissance d'un de ces langages vous permettra de tirer le meilleur profit des exercices. Si CUDA sert aussi aux programmes en Fortran, nous nous limitons ici à CUDA pour C/C++ en utilisant le terme CUDA C. Il s'agit essentiellement de C/C++ permettant d'exécuter des fonctions sur les CPUs et les GPU.

Objectifs d'apprentissage

• Comprendre l'architecture d'un GPU
• Comprendre le déroulement d'un programme CUDA
• Comprendre et gérer les différents types de mémoire GPU
• Écrire et compiler du code CUDA simple avec des exemples

Qu'est-ce qu'un GPU?

Un GPU (pour graphic processing unit), est un processeur monopuce capable d'effectuer des calculs mathématiques rapidement; initialement, on l'utilisait surtout pour produire des rendus d'images. Depuis quelques années, la puissance du GPU sert à accélérer l'exécution de calculs intensifs dans plusieurs domaines de la recherche scientifique de pointe.

Qu'est-ce que CUDA?

CUDA est l'abréviation de Compute Unified Device Architecture. Il s'agit d'un environnement logiciel flexible et d'un modèle de programmation pour le traitement de calculs parallèles intensifs.

Architecture du GPU

Un GPU comporte deux types d'éléments principaux :

• la mémoire globale
  • est semblable à la mémoire du CPU
  • est accessible par un CPU et un GPU
• des multiprocesseurs en continu (SM pour streaming multiprocessor)
  • chaque SM est formé de plusieurs processeurs en continu (SP pour streaming processor)
  • qui effectuent les calculs
  • chaque SM est doté d'une unité de contrôle, de registres, de pipelines d'exécution, etc. qui sui sont propres

Modèle de programmation

Voyons d'abord quelques termes importants :

• Hôte : désigne le CPU et la mémoire principale
• Composant : désigne le GPU et sa mémoire

Le modèle CUDA est un modèle hétérogène qui utilise à la fois le CPU et le GPU. Le code CUDA peut gérer les deux types de mémoires, la mémoire principale du CPU et la mémoire du composant GPU le code exécute aussi les fonctions du GPU appelées noyaux (kernels). Ces fonctions sont exécutées en parallèle par plusieurs fils GPU. Le processus comporte cinq étapes :

1. déclaration et allocation de la mémoire principale et de la mémoire du composant
2. initialisation de la mémoire principale
3. transfert des données de la mémoire principale à la mémoire du composant
4. exécution des fonctions GPU (noyaux)
5. retour des données à la mémoire principale

CUDA Block-Threading Model

CUDA block-threading model where threads are organized into blocks while blocks are further organized into grid.

Given very large number of threads (and in order to achieve massive parallelism one has to use all the threads possible) in CUDA kernel, one needs to organize them somehow. in CUDA, all the threads are structured in threading blocks, the blocks are further organized into grids, as shown on FIg. In dividing the threads we make sure that the following is satisfied:

• threads within a block cooperate via the shared memory
• threads in different blocks can not cooperate

In this model the threads within a block work on the same set of instructions (but perhaps with different data sets) and exchange data between each other via shared memory. Threads in other blocks do the same thing (see Figure).

Threads within a block intercommunicate via shared memory .

Each thread uses IDs to decide what data to work on:

• Block IDs: 1D or 2D (blockIdx.x, blockIdx.y)
• Thread IDs: 1D, 2D, or 3D (threadIdx.x, threadIdx.y, threadIdx.z)

Such model simplifies memory addressing when processing multidimmensional data.

Threads Scheduling

Usually streaming microprocessor (SM) executes one threading block at a time. The code is executed in groups of 32 threads (called Warps). A hardware scheduller is free to assign blocks to any SM at any time. Furthermore, when SM gets the block assigned to it, it does not mean that this particular block will be executed non-stop. In fact, the scheduler can postpone/suspend execution os such block under certain conditions when e.x. data becomes unavailable (indeed, it takes quite some time to read data from the global GPU memory). When it happens, the scheduler takes another threading block which is ready for execution. This is a so called zero-overhead scheduling which makes the execution more stream-lined where SMs are not idling.

GPU Memories in CUDA

There are several type of memories exists for CUDA operations:

• Global memory
  • off-chip, good for I/O, but relatively slow
• Shared memory
  • on-chip, good for thread collaboration, very fast
• Registers& Local Memory
  • thread work space , very fast
• Constant memory

Few Basic CUDA Operations

CUDA Memory Allocation

• cudaMalloc((void**)&array, size)
  • Allocates object in the device memory. Requires address of a pointer of allocated array and size.
• cudaFree(array)
  • Deallocates object from the memory. Requires just a pointer to the array.

CUDA Data Transfers

• cudaMemcpy(array_dest, array_orig, size, direction)
  • Copy the data from either device to host or host to device . Requires pointers to the arrays, size and the direction type (cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice, etc)
• cudaMemcpyAsync
  • Same as cudaMemcpy, but transfers the data asynchronously which means it's not blocking execution of other processes.

First CUDA C Program

The following example shows how to add two numbers on the GPU using CUDA. Note that this is just an exercise, it's very simple, so it will not scale up.

__global__   void add (int *a, int *b, int *c){

	*c = *a + *b;
}
int main(void){
	int a, b, c;
	int *dev_a, *dev_b, *dev_c;
	int size = sizeof(int);

//  allocate device copies of a,b, c
cudaMalloc ( (void**) &dev_a, size);
cudaMalloc ( (void**) &dev_b, size);
cudaMalloc ( (void**) &dev_c, size);

a=2; b=7;
//  copy inputs to device
cudaMemcpy (dev_a, &a, size, cudaMemcpyHostToDevice);
cudaMemcpy (dev_b, &b, size, cudaMemcpyHostToDevice);

// launch add() kernel on GPU, passing parameters
add <<< 1, 1 >>> (dev_a, dev_b, dev_c);

// copy device result back to host
cudaMemcpy (&c, dev_c, size, cudaMemcpyDeviceToHost);

cudaFree ( dev_a ); cudaFree ( dev_b ); cudaFree ( dev_c ); 
}

Are we missing anything ? That code does not look parallel ! Solution: Lets look at what inside the triple brackets in the Kernel call and make some changes :

add <<< N, 1 >>> (dev_a, dev_b, dev_c);

Here we replaced 1 by N, so that N different cuda blocks will be executed at the same time. However, in order to achieve a parallelism we need to make some changes to the Kernel as well:

__global__   void add (int *a, int *b, int *c){

	c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

where blockIdx.x is the unique number identifying a cuda block. This way each cuda block adds a value from a[ ] to b[ ].

CUDA blocks-based parallelism.

Can we again make some modifications in those triple brackets ?

add <<< 1, '''N''' >>> (dev_a, dev_b, dev_c);

Now instead of blocks, the job is distributed across parallel threads. What is the advantage of having parallel threads ? Unlike blocks, threads can communicate between each other ? In other words, we parallelize across multiple threads in the block when massive communication is involved. The chunks of code that can run independently (without much communication) are distributed across parallel blocks.

Advantage of Shared Memory

So far all the memory transfers in the kernel have been done via the regular GPU (global) memory which is relatively slow. Often time we have so many communications between the threads that decreases the performance significantly. In order to address this issue there exist another type of memory called Shared memory which can be used to speed-up the memory operations between the threads. However the trick is that only the threads within a block can communicate. In order to demonstrate the usage of such shared memory we consider the dot product example where two vectors are dot-multipled. Below is the kernel:

__global__   void dot(int *a, int *b, int *c){
        int temp = a[threadIdx.x]*b[threadIdx.x]; 
}

After each thread computed its portion, we need to add everything together. Each threads has to share its data. However, the problem is that each copy of thread's temp is private.This can resolved with the use of shared memory. Below is the kernel with the modifications to account the shared memory usage:

#define N 512
__global__   void dot(int *a, int *b, int *c){
   __shared__ int temp[N];
   temp[threadIdx.x] = a[threadIdx.x]*b[threadIdx.x];
   __syncthreads(); 
   if(threadIdx.x==0){
	int sum; for(int i=0;i<N;i++) sum+= temp[i];
	*c=sum; } 
}

Basic Performance Considerations

Memory Transfers

• PCI-e is extremely slow (4-6 GB/s) compared to both host and device memories
• Minimize Host-to-Device and Device-to-Host memory copies
• Keep data on the device as long as possible
• Sometimes it is not effificient to make the Host (CPU) do non-optimal jobs; executing it on the GPU may still be faster than copying to CPU, executing, and copying back
• Use memcpy times to analyse the execution times

Bandwidth

• Always keep CUDA bandwidth in mind when chaning your code
• Know the theoretical peak bandwidth of the various data links
• Count bytes read/written and compare to the theoretical peak
• Utilize the various memory spaces depending on the situation: global, shared, constant

Common GPU Programming Strategies

• Constant memory also resides in DRAM- much slower access than shared memory
  • BUT, it’s cached !!!
  • highly efficient access for read-only, broadcast
• Carefully divide data acording to access patterns:
  • R Only: constant memory (very fast if in cache)
  • R/W within Block: shared memory (very fast)
  • R/W within Thread: registers (very fast)
  • R/W input/results: global memory (very slow)