Programiranje GPU z CUDA

CUDA (Compute Unified Device Architecture) je paralelna računska arhitektura razvita v NVIDIA in se izvaja v grafičnih karticah (GPU).

Motivacija in prvi koncepti

Razumevanje osnov
Meritve performans

GPU so hitri. Njihova zmogljivost raste hitreje (2x/leto) kot pri normalnih CPU (1.5x/leto). GPU so namenjene za računsko intentivne aplikacije. Glavni namen GPU je prav zaprev za industrijo računalniških igric, vendar se lahko to uporabi tudi za specializirane znanstvene simulacije. Tudi število člankov na tem področju raste. Zmogljivost in performansa GPU se je v zadnjem času primerna za splošno programiranje. Za začetnike se priporočati naslednji dve knjigi:

CUDA by example - bolj didaktična
Programming Massively Parallel Processors - težja za začetnike

Uvod

CUDA je primerna za probleme, ki so zelo probleme ki so zelo paralelni. Se pravi da med deli ni povezave.
Ker je isti program izvajan v vsakem podatkovnem elementu, je pri tem malo zahtev za zahtevne poteke progama.
CUDA tako ni namenjena za vsakovrstne probleme

Programski model

Program v CUDA se sestoji iz ene ali več faz, ki se izvajajo na katerikolem CPU ali na enoti kot je GPU
Koda je napisana v ANSI C razširjena z ukazi
Jedro izvaja veliko število niti, da izrabi praralelizem podatkov
Niti (threads) na GPU so lažje kot pri normalnem CPU
Programska struktura v CUDA:
- Izvajanje začne host
- Ko je funkcija izvedeba se izvajanje prenese v enoto GPU z velikim številom niti
- Vse niti vzpostavljene z jedrom so imenovane grid
Promet na vodilu PCI mora biti minimalen. Za CUDA se podatki prenašajo iz host v napravo. Ti prenosi podatkov so potratni v smislu performanse in jih je potrebno minimirati
- Kompleksnost operacij mora opravičevati ceno prenosa podatkov v GPU.
- Podatki naj se hranijo v GPU čimdlje saj se s tem minimira prenose. To pomeni da naj se vmesni rezultati hranijo na GPU in ne prenašajo na host in nato naza za nadaljnje računanje.
Glavnik kandidati za paralelizacijo z GPU so zanke FOR.

Primer izvajanja jedra v katerem ni vidna FOR zanka, saj le to prevzema funkcija jedra.

// Definicija jedra
__global__ void VecAdd(float *Am float *B, float *B)
{
  int i = threadIdx.x;
  C[i] = A[i] + B[i];
}

int main()
{
  // Izvajanje jedra
  VecAdd<<<1, N>>>(A, B, C);
}

Z threadId vsaka nit izvaja jedro z unikatnim ID, ki je dosegljivo v jedru s spremenljivko threadIdx.. Nit lahko vidimo kot eno ali večdimenzionalno polje. Naslednja koda sešteje dve matriki in shrani rezultat v C.

__global__ void MatAdd(...)
{
  int i = threadIdx.x;
  int j = threadIdx.y;
  C[i][j] = A[i][j] + B[i][j];
}
int main()
{
  dim3 dimBlock(N, N);
  MatAdd<<<1, dimBlock>>>(A, B, C);
}

Dimenzionalnost mreže (grid) se podaja v prvem argumentu sintakse <<< ... >>> v kateri kličemo jedro.

__global__ void MatAdd(float A[N],float B[N], float C[N])
{
   int i = blockIdx.x * blockDim.x + threadIdx.x;
   if (i < N) C[i] = A[i] + B[i]; // Pazi na primer ko je N=1000!!
}
N = 252 * 16;
N = 1000; // alternativa
int main()
{
  
  dim3 dimBlock=256;
  dim3 dimGrid = (N + dimBlock.x - 1) / dimBlock.x ;
  dim3 dimGrid = 16 ; //  alternativa
  MatAdd<<<dimGrid, dimBlock>>>(A, B, C);
}

Dimenzionalnost se lahko prenese še višje na več GPU, kjer imamo tudi blockID in threadId.

Hierarhija pomnilnika

CPU in GPU imata ločene spominske prostore
Host (CPU) upravlja s spominom kartice (GPU)
cudaMalloc(void ** pointerm size_t bytes) - Kličemo iz računalnika, da alociramo globalni spomin.
Kopiranje spomina je možno v vse smeri z:
- cudaMemcpyHostToDevice
- cudaMemcpyDeviceToHost
- cudaMemcpyDeviceToDevice
- cudaMemcpyHostToHost

Primer kopiranja:

int main(void)
{
  float *a_h, *b_h; // poratki v računalniku CPU
  float *a_d, *b_d; // podatki v kartici GPU
  int N = 14, nBytes, i ;
  nBytes = N*sizeof(float);
  a_h = (float *)malloc(nBytes);
  b_h = (float *)malloc(nBytes);
  cudaMalloc((void **) &a_d, nBytes);
  cudaMalloc((void **) &b_d, nBytes);
  for (i=0, i<N; i++)           
     a_h[i] = 100.f + i;
  cudaMemcpy(a_d, a_h, nBytes, cudaMemcpyHostToDevice);
  cudaMemcpy(b_d, a_d, nBytes, cudaMemcpyDeviceToDevice);
  cudaMemcpy(b_h, b_d, nBytes, cudaMemcpyDeviceToHost);
  for (i=0; i< N; i++)       
     assert( a_h[i] == b_h[i] );
  free(a_h); free(b_h); cudaFree(a_d); cudaFree(b_d);
  return 0;
}

Označevanje funkcij

Jedro ima označbo __global__
- Funkcija je klicana iz računalnika in je izvajana v kartici
- Mora vrniti void
Druge CUDA označbe so:
- __device__ - funkcija ki jo kliče kartica in je v kartici
- __host__ - privzeta funkcija na računalniku

Označevanje spremenljivk

__device__ - shranjeno v globalnem spominu GPU in dosegljivo za vse niti
__constant__ - samo za branje v GPU
__shared__ - Hiter pomnilnik za rabo v jedru. Dosegljiv iz vseh niti.

Poganjanje jeder

Sintaksa kernel<<<dim3 dG, dim3 dB>>>(...) mogoča naslednje konfiguracije:dG - dimension and size of grid in blocks

dG - dimenzije grid in blokov
- Dimenzija v x in y
- bloki se poganjajo na gridu: dG.x * dG.y
- hardverska omejitev 65,535 blokov na dimenzijo
dB - dimenzija in velkost blokov v nitih
- Tri-dimenzionalno: x, y, and z. zaenkrat samo 2D x in y čeprav jezik predvideva tudi z za prihodnje arhitekture
- Število niti v bloku je: dB.x * dB.y * dB.z
- Primeri: (512,1,1), (8,16,2) or (16,16,2)
- Ni dovoljeno: (32,32,1)
- hardverska omejitev 65,535 blokov na dimenzijo

Prvi program

#include <stdio.h>
#include <cuda.h>
// Jedro ki ga izvaja CUDA kartica
__global__ void square_array(float *a, int N)
{
}

// glavna rutina, ki se izvaja na računalniku (host)
int main(void)
{
   // Kazalec na host in polja v kartici
   // Število elementov v poljih
   // Alokacija polj na hostu
   // Alokacija polj na kartici
   // Inicializacija polja na hostu in kopiranje v CUDA kartico
   //  Računaj na kartici
   // Poberi rezultate iz kartice in shrani v računalnik
   // Izpiši rezultate
   // Čiščenje
}

Posamezni deli programa so:

Jedro ki ga izvaja CUDA kartica

__global__ void square_array(float *a, int N)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx<N) a[idx] = a[idx] * a[idx];
}

Kazalec na polja v hostu in v kartici

int main(void)
{
   float *a_h, *a_d; // Kazalca na polja v host & device
   const int N = 10; // Število elementov v polju
.
.
.
}

Alokacija spomina

size_t size = N * sizeof(float);
  a_h = (float *)malloc(size); // Alociraj polje na hostu
  cudaMalloc((void **) &a_d, size); // alociraj polje na kaAllocate array on device

Inicializacija polja na hostu in kopiranje v CUDA kartico

for (int i=0; i<N; i++) a_h[i] = (float)i;
    cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);

Računaj na kartici

int block_size = 4;
   int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);
   square_array <<< n_blocks, block_size >>> (a_d, N);

Poberi rezultate iz kartice in shrani v računalnik

cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost);

Izpiši rezultate in počisti spomin

for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
   // Čiščenje
   free(a_h); cudaFree(a_d);

Prevajanje in poganjanje programa a.cu

$ nvcc a.cu -o a
$ ./a
0 0.000000
1 1.000000
2 4.000000
3 9.000000
4 16.000000
5 25.000000
6 36.000000
7 49.000000
8 64.000000
9 81.000000

Hardverske secifikacije

Naslednji program izpiše lastnosti CUDA kartice:

#include <stdio.h>
#include <cuda.h>

int main(void)
{
 cudaDeviceProp prop;
 int count;
 cudaGetDeviceCount( &count );
 for (int i=0; i< count; i++) {
  cudaGetDeviceProperties( &prop, i ) ;
  printf( " --- General Information for device %d ---\n", i );
  printf( "Name: %s\n", prop.name );
  printf( "Compute capability: %d.%d\n", prop.major, prop.minor );
  printf( "Clock rate: %d\n", prop.clockRate );
  printf( "Device copy overlap: " );
  if (prop.deviceOverlap)
    printf( "Enabled\n" );
  else
    printf( "Disabled\n");
  printf( "Kernel execution timeout : " );
  if (prop.kernelExecTimeoutEnabled)
    printf( "Enabled\n" );
  else
    printf( "Disabled\n" );
  printf( " --- Memory Information for device %d ---\n", i );
  printf( "Total global mem: %ld\n", prop.totalGlobalMem );
  printf( "Total constant Mem: %ld\n", prop.totalConstMem );
  printf( "Max mem pitch: %ld\n", prop.memPitch );
  printf( "Texture Alignment: %ld\n", prop.textureAlignment );
  printf( " --- MP Information for device %d ---\n", i );
  printf( "Multiprocessor count: %d\n",
  prop.multiProcessorCount );
  printf( "Shared mem per mp: %ld\n", prop.sharedMemPerBlock );
  printf( "Registers per mp: %d\n", prop.regsPerBlock );
  printf( "Threads in warp: %d\n", prop.warpSize );
  printf( "Max threads per block: %d\n", prop.maxThreadsPerBlock );
  printf( "Max thread dimensions: (%d, %d, %d)\n",
                prop.maxThreadsDim[0], prop.maxThreadsDim[1],
  prop.maxThreadsDim[2] );
  printf( "Max grid dimensions: (%d, %d, %d)\n",
  prop.maxGridSize[0], prop.maxGridSize[1],
  prop.maxGridSize[2] );
  printf( "\n" );
         }
}

Na kartici, kjer smo izvajali teste je program izpisal:

--- General Information for device 0 ---
Name: GeForce GTX 295
Compute capability: 1.3
Clock rate: 1242000
Device copy overlap: Enabled
Kernel execution timeout : Disabled
 --- Memory Information for device 0 ---
Total global mem: 939261952
Total constant Mem: 65536
Max mem pitch: 262144
Texture Alignment: 256
 --- MP Information for device 0 ---
Multiprocessor count: 30
Shared mem per mp: 16384
Registers per mp: 16384
Threads in warp: 32
Max threads per block: 512
Max thread dimensions: (512, 512, 64)
Max grid dimensions: (65535, 65535, 1)

 --- General Information for device 1 ---
Name: GeForce GTX 295
Compute capability: 1.3
Clock rate: 1242000
Device copy overlap: Enabled
Kernel execution timeout : Disabled
 --- Memory Information for device 1 ---
Total global mem: 939261952
Total constant Mem: 65536
Max mem pitch: 262144
Texture Alignment: 256
 --- MP Information for device 1 ---
Multiprocessor count: 30
Shared mem per mp: 16384
Registers per mp: 16384
Threads in warp: 32
Max threads per block: 512
Max thread dimensions: (512, 512, 64)
Max grid dimensions: (65535, 65535, 1)

 --- General Information for device 2 ---
Name: GeForce GTX 295
Compute capability: 1.3
Clock rate: 1242000
Device copy overlap: Enabled
Kernel execution timeout : Disabled
 --- Memory Information for device 2 ---
Total global mem: 939261952
Total constant Mem: 65536
Max mem pitch: 262144
Texture Alignment: 256
 --- MP Information for device 2 ---
Multiprocessor count: 30
Shared mem per mp: 16384
Registers per mp: 16384
Threads in warp: 32
Max threads per block: 512
Max thread dimensions: (512, 512, 64)
Max grid dimensions: (65535, 65535, 1)

 --- General Information for device 3 ---
Name: GeForce GTX 295
Compute capability: 1.3
Clock rate: 1242000
Device copy overlap: Enabled
Kernel execution timeout : Disabled
 --- Memory Information for device 3 ---
Total global mem: 939261952
Total constant Mem: 65536
Max mem pitch: 262144
Texture Alignment: 256
 --- MP Information for device 3 ---
Multiprocessor count: 30
Shared mem per mp: 16384
Registers per mp: 16384
Threads in warp: 32
Max threads per block: 512
Max thread dimensions: (512, 512, 64)
Max grid dimensions: (65535, 65535, 1)

Deljeni spomin

Je hitrejši kot lokalni in globalni spomin(100x). Deljeni spomin deklariramo z __shared__. Velikost bloka v deljenem spominu je omejena na 16x16 oz TIME_DIM=16.

Primerjava z in brez deljenega spomina

__global__ void coalescedMultiply(float *a, float* b, float *c, int N)
{
  __shared__ float aTile[TILE_DIM][TILE_DIM];
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
  for (int i = 0; i < TILE_DIM; i++) {
    sum += aTile[threadIdx.y][i]* b[i*N+col];
  }
  c[row*N+col] = sum;
}
// Brez deljenega spomina
__global__ void simpleMultiply(float *a, float* b, float *c, int N)
{
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  for (int i = 0; i < TILE_DIM; i++) {
    sum += a[row*TILE_DIM+i] * b[i*N+col];
  }
  c[row*N+col] = sum;
}

Prijavi se za komentar

Main menu

High Performance Computing

Navigacija

Nahajate se tukaj