Dezvoltare de cod pentru CUDA

Instalare CUDA framework

Are mai multe etape:

  1. instalarea unui driver pentru placa video cu suport pentru CUDA. In Linux se poate folosi packet manager (in functie de distributie: apt, yum, yast etc.) sau se poate rula scriptul oferit de nVidia. Desigur, pentru rularea scriptului este necesara oprirea lui X11 si pornirea sistemului de operare intr-un run-level de cel mult 3.
  2. instalarea de toolkit: este simpla atat in Linux, unde trebuie rulat scriptul descarcat odata cu toolkit-ul si urmati pasii acestuia, cat si in Windows, unde avem la dispozitie un installer.
  3. optional instalarea GPU Computing SDK, cu mostre program, sub forma de proiecte Visual Studio pentru versiunea de Windows sau cu Makefile pentru versiunea Linux.

Toate componentele sunt disponibile atat pentru arhitecturi 32-bit, cat si 64-bit. La resurse, puteti gasi un tutorial pentru instalarea frameworkului.

Compilare

Codul consta din doua componente:

  • componenta host se refera la codul ce va rula pe CPU;
  • componenta device se refera la codul ce va rula pe GPU (alcatuita din kernel-uri).

Compilatorul oferit de nVidia este nvcc, el compiland codul ce va rula pe device si apeland un compilator clasic (e.g. gcc pe Linux) pentru codul host. La final apeleaza linkerul si construieste executabilul final.

Setul de instructiuni PTX reprezinta echivalentul limbajului de asamblare. Programele pot fi scrise direct in cod PTX, pentru control maxim asupra GPU. In timpul procesului de compilare, codul C al kernelului este transformat in cod PTX, acesta putand fi inspectat daca se foloseste flagul –keep pentru nvcc.

 nvcc gpu_code.c -o gpu_code 

Flaguri utile la compilarea cu NVCC:
-keep: pastreaza toate fisierele intermediare, util pentru inspectarea transformarilor si optimizarilor executate de compilator
-arch=sm_xx: seteaza compute capability (sm_10 .. sm_13)
–deviceemu: Compileaza programul pentru rularea in emulator, in locul executiei pe device
-g -G: Debug mode

Înainte de a putea folosi compilatorul nVidia, trebuie să încarcaţi modulul corespunzător:

module load libraries/cuda-5.0 

Compute capability

Pentru a suporta diferitele chipuri nVidia, CUDA permite specificarea capabilitatilor platformei tinta la compilare:

 nvcc -arch=sm_## gpu_code.c

De exemplu, se foloseste sm_13 pentru compute capability 1.3

Implicit, valoarea pentru compute_capability este 1.0. Acest lucru poate fi important in anumite cazuri, spre exemplu suportul pentru operatii in virgula mobila dubla precizie este prezent in compute_capability 1.3. In versiunile precedente, compilatorul transforma automat declaratiile variabilelor din kernel din double in float.

O lista a placilor video si compute_capability aferent poate fi gasită in CUDA Programming Guide, Appendix A. Pentru a afla compute_capability precum și alți parametri importanți ai placii grafice (cantitatea de shared memory, numărul de blocuri etc), se poate rula aplicația deviceQuery din Nvidia GPU Computing SDK (aflat pe in directorul: /opt/tools/libraries/cuda-5.0/samples/1_Utilities, trebuie încărcat în prealabil modulul pentru CUDA). Tot în acest director aveți și o serie de alte exemple relevante.

Un exemplu de output al deviceQuery:

Device 0: "GeForce GTX 260"
  CUDA Driver Version:                           2.30
  CUDA Runtime Version:                          2.30
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         3
  Total amount of global memory:                 938803200 bytes
  Number of multiprocessors:                     24
  Number of cores:                               192
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 16384
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          262144 bytes
  Texture alignment:                             256 bytes
  Clock rate:                                    1.30 GHz
  Concurrent copy and execution:                 Yes
  Run time limit on kernels:                     Yes
  Integrated:                                    No
  Support host page-locked memory mapping:       Yes
  Compute mode:                                  Default (multiple host threads can use this device simultaneously)

Alocarea memoriei

Memoria pentru device se aloca folosind functia cudaMalloc(). Exista rutine speciale pentru alocarea de array-uri 2D si 3D. Pentru alocarea de memorie unidimensionala:

cudaError_t  cudaMalloc(void** devPtr, size_t size) 			
 
    	devPtr 	- Pointer spre memoria device alocata
    	size 	- Cantitatea de memorie ceruta

Eliberarea memoriei se face cu functia cudaFree().

Copierea datelor in memoria video

Datele pot fi mutate intre memoria host si memoria device in mod sincron sau asincron. Pentru performanta maxima, se recomanda folosirea modului asincron.

Exista rutine speciale pentru copierea array-urilor 2D sau 3D. Pentru copierea sincrona a unui array unidimensional:

cudaError_t  cudaMemcpy (void * dst, const void * src, size_t count, enum cudaMemcpyKind kind) 
 
    	dst 	- Destination memory address
    	src 	- Source memory address
    	count 	- Size in bytes to copy
    	kind 	- Tipul transferului: cudaMemcpyHostToHost, cudaMemcpyHostToDevice, 
                  cudaMemcpyDeviceToHost sau cudaMemcpyDeviceToDevice

Model de utilizare:

  • cudaMemcpy(…, cudaMemcpyHostToDevice) - copierea datelor (input) in memoria device-ului
  • rulare kernel
  • cudaMemcpy(…, cudaMemcpyDeviceToHost) - copierea datelor (output) din memoria device-ului

Rularea unui kernel

La rularea unei functii kernel se specifica o configuratie pentru executie: sunt dati drept parametri configuratia blocurilor de threaduri si configuratia threadurilor intr-un bloc.

__global__ void KernelFunc(args...){..}
 
dim3   DimGrid(100, 50);    // 5000 thread blocks 
dim3   DimBlock(4, 8, 8);   // 256 threads per block 
 
KernelFunc<<< DimGrid, DimBlock >>>(args...);

In exemplul de mai sus este definita o configuratie bidimensionala de blocuri si o configuratie tridimensionala de threaduri. Un bloc de threaduri poate avea una, doua sau trei dimensiuni. O grila de blocuri poate avea una sau doua dimensiuni.

Un kernel poate fi apelat fara folosirea de variabile intermediare pentru definirea configuratiei threadurilor:

 KernelFunc<<< 32, 16 >>>(args...);		// grila unidimensionala de 32 blocuri, 16 threaduri per bloc 

Apelurile de kernel sunt operatii asincrone. Este nevoie de o sincronizare explicita pentru a astepta terminarea executiei.

Identificarea unui thread

Fiecare thread se poate identifica cu pozitia in cadrul blocului (threadIdx) si pozitia blocului in cadrul grilei (blockIdx). Astfel se pot stabili operatiile ce urmeaza a fi executate de fiecare thread.

Variabilele

threadIdx={x,y,z} 

si

blockIdx={x,y} 

sunt accesibile oricarui CUDA thread. Un exemplu de utilizare in adunarea elementelor a doua matrici prezentata in sectiunile urmatoare.

Cooperare si sincronizare

Threadurile dintr-un bloc pot coopera schimband date prin intermediul shared memory. Variabilele alocate in shared memory au prefixul __shared__.

 __shared__ float A[BLOCK_SIZE][BLOCK_SIZE]; 

Sincronizarea se face doar pentru threadurile din acelasi bloc,. __syncthreads() actioneaza ca o bariera, executia fiecarui thread continuand doar dupa ce fiecare a intrat in bariera.

Considerente pentru dezvoltarea de programe pentru CUDA

  • folosirea a cat mai multe threaduri
  • limitarea accesului la memoria globala
  • folosirea memoriei shared pe post de cache programabil:
    • datele sunt partitionate in segmente ce incap in memoria shared
    • executa calculele pe segmentul din memoria shared
    • copiaza rezultatele din memoria shared in memoria globala
  • folosirea memoriei constante pentru datele read only
  • evitarea ramurilor divergente (divergent branches) si a serializarii warpurilor (warp serialization)
  • preferarea calcularii on-the-fly in loc de precalculare+incarcare din memoria globala

Exemple

Dublarea valorilor dintr-un array

__global__ void doubleArrayOnDevice(float* temp, int size)
{
	int index = blockIdx.x*blockDim.x + threadIdx.x;
	if (index<N) temp[index] = 2*temp[index];
}

Cod complet: doublearr.cu.tar

Adunare de matrice

matadd.cu
// Kernel definition
__global__ void MatAdd(float* A, float* B, float* C)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  int j = blockIdx.y * blockDim.y + threadIdx.y;
 
  if (i < N && j < N)
     C[i*N+j] = A[i*N+j] + B[i*N+j];
}
 
int main()
{
 
  float a[N*N], b[N*N], c[N*N]; 	// host matrices
  float *A, *B, *C;			// device matrices
 
  ...
 
  // Kernel invocation
  dim3 dimBlock(BLK, BLK);
  dim3 dimGrid((N + dimBlock.x - 1) / dimBlock.x, (N + dimBlock.y - 1) / dimBlock.y);
  MatAdd<<<dimGrid, dimBlock>>>(A, B, C);
 
  ...
 
}

Inmultire de matrice

Memoria shared este mult mai rapida decat memoria globala. Deci orice oportunitate de a inlocui accesul la memoria globala cu cu acces la cea shared trebuie exploatata. Urmatorul exemplu de multiplicare de matrice ilustreaza acest lucru.

Inmultire directa

Codul urmator este o implementare directa a multiplicarii matricilor si nu foloseste memoria shared.

Fiecare thread citeste o linie din A, o coloana din B si calculeaza elementul corespondent din C. A este deci citita din memoria globala de B.width ori iar B este citit de A.height ori.

matmul_noshare.cu
// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
  int width;
  int height;
  float* elements;
} Matrix;
 
// Thread block size
#define BLOCK_SIZE 16
 
// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
 
// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
  // Load A and B to device memory
  Matrix d_A;
 
  d_A.width = A.width; 
  d_A.height = A.height;
 
  size_t size = A.width * A.height * sizeof(float);
 
  // TODO: alocati matricea d_A pe device si puneti in ea date din matricea A de pe host
 
  Matrix d_B;
 
  d_B.width = B.width; 
  d_B.height = B.height;
 
  size = B.width * B.height * sizeof(float);
 
  // TODO: alocati matricea d_B pe device si puneti in ea date din matricea B de pe host
 
  // Allocate C in device memory
  Matrix d_C;
 
  d_C.width = C.width; 
  d_C.height = C.height;
 
  size = C.width * C.height * sizeof(float);
  cudaMalloc((void**)&d_C.elements, size);
 
  // Invoke kernel
  dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
 
  MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
 
  // Read C from device memory
  // TODO: copiati in matricea C de pe host rezultatul din matricea d_C de pe device
 
  // Free device memory
  cudaFree(d_A.elements);
  cudaFree(d_B.elements);
  cudaFree(d_C.elements);
}
 
  // Matrix multiplication kernel called by MatrixMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
  // Each thread computes one element of C
  // by accumulating results into Cvalue
  float Cvalue = 0;
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
 
  for (int e = 0; e < A.width; ++e)
     Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
 
  C.elements[row * C.width + col] = Cvalue;
}

Inmultire de matrice optimizata pentru CUDA

Codul urmator este o implementare pentru inmultire de matrice ce profita de memoria shared. In aceasta varianta, fiecare bloc de threaduri este responsabil pentru calcului unei singure sub-matrice Csub a lui C iar fiecare thread din bloc este responsabil cu calculul unui singur element din Csub.

Csub este egal cu produsul a doua matrice: o submatrice a lui A de dimensiune (A.width, block_size), care are aceeasi indici de linie ca si Csub, si submatricea B de dimensiune (block_size, A.width) ce are aceeasi indici de coloana ca si Csub.

Pentru a incapea in resursele device-ului, cele doua matrice sunt impartite in atatea submatrici patratice de dimensiune block_size de cat este nevoie. Csub este calculata ca suma produselor acestor matrice. Fiecare din aceste produse este obtinut prin incarcarea celor doua matrice corespondente din memoria globala in cea shared, cu un thread incarcand un singur element al matricei.

Apoi fiecare thread calculeaza un element al produsului. Fiecare thread acumuleaza rezultatul fiecaruia din aceste produse intr-un registru si, la final, scrie rezultatul in memoria globala.

matmul_share.cu
// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
 
typedef struct {
  int width;
  int height;
  int stride;
  float* elements;
} Matrix;
 
// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{
  return A.elements[row * A.stride + col];
}
 
// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col, float value)
{
  A.elements[row * A.stride + col] = value;
}
 
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A
__device__ Matrix GetSubMatrix(Matrix A, int row, int col)
{
  Matrix Asub;
  Asub.width = BLOCK_SIZE;
  Asub.height = BLOCK_SIZE;
  Asub.stride = A.stride;
  Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col];
 
  return Asub;
}
 
// Thread block size
#define BLOCK_SIZE 16
 
// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
 
// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix N, Matrix C)
{
  // Load A and B to device memory
  Matrix d_A;
 
  d_A.width = d_A.stride = A.width; 
  d_A.height = A.height;
 
  size_t size = A.width * A.height * sizeof(float);
 
  // TODO: alocati matricea d_A pe device si puneti in ea date din matricea A de pe host
 
  Matrix d_B;
 
  d_B.width = d_B.stride = B.width; 
  d_B.height = B.height;
 
  size = B.width * B.height * sizeof(float);
 
  // TODO: alocati matricea d_B pe device si puneti in ea date din matricea B de pe host
 
  // Allocate C in device memory
  Matrix d_C;
 
  d_C.width = d_C.stride = C.width; 
  d_C.height = C.height;
 
  size = C.width * C.height * sizeof(float);
 
  cudaMalloc((void**)&d_C.elements, size);
 
  // Invoke kernel
  dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
 
  MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
 
 
  // Read C from device memory
  // TODO: copiati in matricea C de pe host rezultatul din matricea d_C de pe device
 
  // Free device memory
  cudaFree(d_A.elements);
  cudaFree(d_B.elements);
  cudaFree(d_C.elements);
}
 
// Matrix multiplication kernel called by MatrixMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
  // Block row and column
  int blockRow = blockIdx.y;
  int blockCol = blockIdx.x;
 
  // Each thread block computes one sub-matrix Csub of C
  Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
 
  // Each thread computes one element of Csub
  // by accumulating results into Cvalue
  float Cvalue = 0;
 
  // Thread row and column within Csub
  int row = threadIdx.y;
  int col = threadIdx.x;
 
  // Loop over all the sub-matrices of A and B that are
  // required to compute Csub
  // Multiply each pair of sub-matrices together
  // and accumulate the results
 
  for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) 
  {
   // Get sub-matrix Asub of A
   Matrix Asub = GetSubMatrix(A, blockRow, m);
 
  // Get sub-matrix Bsub of B
  Matrix Bsub = GetSubMatrix(B, m, blockCol);
 
  // Shared memory used to store Asub and Bsub respectively
  __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
  __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
 
  // Load Asub and Bsub from device memory to shared memory
  // Each thread loads one element of each sub-matrix
  As[row][col] = GetElement(Asub, row, col);
  Bs[row][col] = GetElement(Bsub, row, col);
 
  // Synchronize to make sure the sub-matrices are loaded
  // before starting the computation
  __syncthreads();
 
  // Multiply Asub and Bsub together
  for (int e = 0; e < BLOCK_SIZE; ++e)
    // TODO: calculati un element folosind datele din memoria shared
 
  // Synchronize to make sure that the preceding
  // computation is done before loading two new
  // sub-matrices of A and B in the next iteration
  __syncthreads();
}
 
  // Write Csub to device memory
  // Each thread writes one element
  SetElement(Csub, row, col, Cvalue);
}
asc/gpucookbook/dezvoltare.txt · Last modified: 2014/05/09 13:22 by dan.dragomir
CC Attribution-Share Alike 3.0 Unported
www.chimeric.de Valid CSS Driven by DokuWiki do yourself a favour and use a real browser - get firefox!! Recent changes RSS feed Valid XHTML 1.0