CUDA »

CUDA

Podstrony


Wszystkie notatki na jednej stronie

CUDA API

Do dyspozycji są dwa interfejsy programisty:

  • CUDA C Runtime API: funkcje cuda*()
    • zaimplementowana w dynamicznie ładowanej bibliotece cudart
    • wsparcie dla Open GL? i Direct 3 D?
  • Driver API: funkcje cu*()
    • większa kontrola nad urządzeniami, np. jeden wątek CPU może obsługiwać wiele GPU.
    • kompilacja PTX Just-In-Time
    • możliwość kompilowania kodu hosta dowolnym kompilatorem (nie ma C extensions)
    • brakuje emulatora
    • większa złożoność kodu

Kod wykonywany na GPU nie zależy od użytego API.

Migracja z Runtime API do Driver API jest fatwa, ale odwrotnie może być problem.

Zarządzanie urządzenami

Funkcje wywoływane przez CPU umożliwiające zarządzanie GPU:

  • cudaGetDeviceCount( int* count )
  • cudaSetDevice( int device )
  • cudaGetDevice( int *current_device )
  • cudaGetDeviceProperties( cudaDeviceProp* prop, int device )
  • cudaChooseDevice( int *device, cudaDeviceProp* prop )

Domyślnym urządzeniem GPU jest device 0.

Jeden wątek CPU może kontrolować jedno GPU. Wielordzeniowej CPU może kontrolować jedno urządzenie GPU przez serializację odwołać dokonywaną przez sterownik.

Obsługa CUDA Error

Funkcje umożliwiające interpretację wyników wywołania wszystkich funkcji CUDA:

  • cudaError_t cudaGetLastError(void)
  • char* cudaGetErrorString(cudaError_t code)

Makro HANDLE_ERROR

Makro do przechwytywania błędów autorstwa Jasona Sandersa 1:

static void HandleError( cudaError_t err, const char *file,  int line ) {
        if (err != cudaSuccess) {
                printf( "%s in %s at line %d\n", cudaGetErrorString( err ),  file, line );
                exit( EXIT_FAILURE );
        }
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
1 “CUDA By Example An Introduction to General-Purpose GPU Programming” by Jason Sanders and Edward Kandrot, Addison-Wesley, Upper Saddle River, NJ, 2011

Pomiar wydajności

Podstawowym zastosowaniem zdarzeń w CUDA jest pomiar wydajności :)

Przykład kodu:

unsigned int timer = 0;

  CUT_SAFE_CALL(cutCreateTimer(&timer));
  CUT_SAFE_CALL(cutStartTimer(timer));

    // obliczenia na GPU
    kernel<<< grid, threads >>>((int*) d_data);

  cudaThreadSynchronize();  // czekamy na koniec wszelkich obliczeń

  CUT_SAFE_CALL( cutStopTimer( timer));
    printf( "\nGPU Processing time: %f (ms)\n\n", cutGetTimerValue( timer));
  CUT_SAFE_CALL( cutDeleteTimer( timer));

Inny wariant (zaczerpnięty z 2):

cudaEvent_t start, stop;
HANDLE_ERROR( cudaEventCreate( &start ) );
HANDLE_ERROR( cudaEventCreate( &stop ) );
HANDLE_ERROR( cudaEventRecord( start, 0 ) );

...

HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
HANDLE_ERROR( cudaEventSynchronize( stop ) );
  float elapsedTime;
HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime, start, stop ) );
  printf( "Time to generate: %3.1f ms\n", elapsedTime );
HANDLE_ERROR( cudaEventDestroy( start ) );
HANDLE_ERROR( cudaEventDestroy( stop ) );

...
1 “CUDA By Example An Introduction to General-Purpose GPU Programming” by Jason Sanders and Edward Kandrot, Addison-Wesley, Upper Saddle River, NJ, 2011

2 “CUDA By Example An Introduction to General-Purpose GPU Programming” by Jason Sanders and Edward Kandrot, Addison-Wesley, Upper Saddle River, NJ, 2011

Rodzaje pamięci

  • Register - rejestry on-chip threadu,
    • dostępne jedynie w kernelu,
    • nieadresowalna, do przechowywania zmiennych
  • Local memory - pamięć off-chip dla threadu, zazwyczaj w DRAM
    • dostępna na czas wykonywania kernelu,
  • Shared memory - pamięć on-chip wspólna dla bloku thread'ów
    • 16KB = 16 banków po 32b słowa
    • pamięć wspólna dla wszystkich wątków w bloku (R/W)
    • szybka jak rejestry jeśli nie występują konflikty dostępu
    • dostępna tak długo, jak długo wykonywany jest blok
  • Global (device) memory
    • dostępna dla wszystkich thread'ów (GPU) oraz dla hosta (CPU)
    • dostępna od momentu allokacji do deallokacji
  • Constant
  • Texture

Shared Memory

Użycie pamięci wspólnej:

  • o znanym rozmiarze przy kompilacji
__global__ void kernel(...) {
  ...
    __shared__ float sData[256];
  ...
}

int main(void) {
...
    kernel<<<nBlocks,blockSize>>>(...);
...
}
  • rozmiar znany w momencie wywołania kernela:
__global__ void kernel(...) {
...
   extern __shared__ float sData[];
...
}

int main(void) {
...
   smBytes = blockSize*sizeof(float);
   kernel<<<nBlocks, blockSize, smBytes>>>(...);
...
}

Dostęp do pamięci

  • Każdy wątek może:
    • R/W per-thread registers
    • R/W per-thread local memory
    • R/W per-block shared memory
    • R/W per-grid global memory
    • Read only per-grid constant memory
    • Read only per-grid texture memory
  • Program hosta (CPU) ma jedynie dostęp R/W do:
    • global memory
    • constant memory
    • texture memory

Rys. 1. Możliwe dostępy do pamięci (rodzina G80).

Zarządzanie pamięcią

Podstawowe funkcje:

  • cudaMalloc(void **pointer, size_t nbytes)
  • cudaMemset(void *pointer, int value, size_t count)
  • cudaFree(void *pointer)

Przykład:

int n = 1024;
int nbytes = 1024*sizeof(int);
int *d_a = 0;
cudaMalloc( (void**)&d_a, nbytes );
cudaMemset( d_a, 0, nbytes);
cudaFree(d_a);

Kopiowanie danych

 cudaMemcpy(void *dst, void *src, size_t nbytes, enum cudaMemcpyKind direction);
  • enum cudaMemcpyKind
    • cudaMemcpyHostToDevice
    • cudaMemcpyDeviceToHost
    • cudaMemcpyDeviceToDevice

Pamięć page-locked po stronie host'a

W celu przyśpieszenia transferu danych między hostem a kartą (CPU<->GPU), wykorzystuje się allokowanie pamięci w specjalnym trybie page-locked. System operacyjny gwarantuje, że taka pamięć zawsze będzie dostępna w pamięci RAM, nigdy nie zostanie odłożona np. na dysk. Pamięć taką allokuje się po stronie hosta i używa do kopiowania danych do/z karty graficznej. Operacje kopiowania są bardzo wydajne. Dodatkowo, w nowszym sprzęcie (minimum CUDA 2.2) możliwy jest dostęp z poziomu kernela do pamięci page-locked hosta. Allokowanie dużych obszarów w trybie page-locked może wpływać niekorzystnie na stabilność i wydajność całego systemu operacyjnego!

Do zarządzania taką pamięci służą dwie funkcje: cudaMallocHost() oraz cudaFreeHost().

tId - rozróżnianie wątków

Identyfikator wątku (?) obliczany jest z indeksu wątku oraz rozmiarów i indeksu bloku:

int tid = threadIdx.x + blockIdx.x * blockDim.x;

Przykład użycia do adresowania tablic:

__global__ void add( int *a, int *b, int *c ) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while (tid < N) {
       c[tid] = a[tid] + b[tid];
       tid += blockDim.x * gridDim.x;
   }
}

Adresowanie obrazu

Przy użyciu identyfikatorów wątku może zaadresować dwuwymiarowe dane, np. wsp. piksela oraz jego offset w tablicy danych:

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;

Pamięć shared memory jest szybka jak rejestry, pod warunkiem, że nie występują konflikty w dostępnie.

Cytad z dokumentacji:

  • The fast cases:
    • All 16 threads of a half-warp access different banks: no bank conflict
    • All 16 threads of a half-warp access the same address: broadcast
  • The slow case:
    • Multiple threads in the same half-warp access different values in the same bank
    • Must serialize the accesses
    • Cost = max # of values requested from one of the 16 banks

Technika Padding umożliwia redukcję konfliktów pamięci współdzielonej multiprocesora (Shared memory bank conflicts). Polega na modyfikacji indeksów:

  • po wyliczeniu adresu, np.:
    adres = 2* stride * thid;
  • zmieniamy na:
    adres += (adres / 16); // dzielimy przez NUM_BANKS 

Taka sztuczka rozwiązuje większość problemów :)

Przykład padding'u bez dzielenia (na odstawie dokumentacji Davida Kirk'a):

const int LOG_NUM_BANKS = 4; // 16 banks on G80
int tid = threadIdx.x;
int s = 1;
// Traversal from leaves up to root
for (d = n>>1; d > 0; d >>= 1)
{
        if (thid <= d)
        {
        int a = s*(2*tid); int b = s*(2*tid+1)
        a += (a >> LOG_NUM_BANKS); // insert pad word
        b += (b >> LOG_NUM_BANKS); // insert pad word
             shared[a] += shared[b];
   }
}
 

Zmienne wbudowane

dim3 gridDim
Wymiary siatki w blokach (gridDim.z nieużywana)
dim3 blockDim
Wymiar bloku wątków
dim3 blockIdx
Indeks bloku w siatce (grid), 1D lub 2D
dim3 threadIdx
Indeks wątku w bloku, 1D-2D-3D

Organizacja w grid/block/thred dla danych 2D

Użycie

Zmienne wbudowane używane są przede wszystkim podczas wywoływania kernela CUDA, np.:

kernel<<<dim3 grid, dim3 block>>>(...);

Wbudowane funkcje matematyczne

Każdy multiprocesor zawiera zawiera specializowane układy do sprzętowego obliczania podstawowych funkcji matematycznych. Funkcje te obliczane są na liczbach zmiennoprzecinkowych pojedynczej precyzji (float) i mają ograniczoną dokładność. Za to wykonują się bardzo szybko.

Wybrane sprzętowe funkcje (pełna lista znajduje się w dokumentacji):

Operacja lub funkcjaSprzętowo obliczana funkcja na GPU
x/y__fdividef(x,y)
sinf(x)__sinf(x)
cosf(x)__cosf(x)
tanf(x)__tanf(x)
sincosf(x,sptr,cptr)__sincosf(x,sptr,cptr)
logf(x)__logf(x)
log2f(x)__log2f(x)
log10f(x)__log10f(x)
expf(x)__expf(x)
exp10f(x)__exp10f(x)

Koszty obliczeniowe

Do obliczenia jednej instrukcji dla warp'a. multiprocesor potrzebuje (G80, na podstawie CUDA Guide 1v1):

  • 4 taktów dla: floating-point add, floating-point multiply, floating-point, multiply-add, integer add, bitwise operations, compare, min, max, type conversion instruction
  • 16 taktów dla: odwrotności, odwrotności pierwiastka kwadratowego i funkcji wbudowanych typu __log(x)
  • 32 taktów dla: __sin(x), __cos(x), __exp(x)

Dzielenie całkowitoliczbowe oraz operacja modulo są kosztowne obliczeniowo! Należy je zastępować operacjami bitowymi, no. dla n będącego wielokrotnością dwójki:

  • (i/n) odpowiada (i>>log2(n))
  • (i%n) odpowiada (i$(n-1))

Szybkie obliczenie indeksu

Pomnożenie dwóch liczb 32b całkowitych (int) trwa 16cykli zegara GPU, natomiast w użycie __mul24 kozsztuje jedynie 4 cykle. Przykład użycia:

long idx = __mul24( blockIdx.x , blockDim.x ) + threadIdx.x; // 4 cykle na mnożenie!

HomePage

Dydaktyka

Materiały do pobrania

Praca naukowa

Programowanie

Projekty

Lista stron


edit

Page last modified on 11.06.2011, 17:37 - Powered by PmWiki

|