From Mirek Socha zapraszam!

CUDA: CUDA

Podstrony

CUDA /
API
CUDA
Device
Error
Events
Host
Math
Memory
NVCC
Padding
RecentChanges
Streams
Tid
Variables

Wszystkie notatki na jednej stronie

CUDA API

Do dyspozycji są dwa interfejsy programisty:

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:

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:

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

Shared Memory

Użycie pamięci wspólnej:

__global__ void kernel(...) {
  ...
    __shared__ float sData[256];
  ...
}

int main(void) {
...
    kernel<<<nBlocks,blockSize>>>(...);
...
}
__global__ void kernel(...) {
...
   extern __shared__ float sData[];
...
}

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

Dostęp do pamięci

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

Zarządzanie pamięcią

Podstawowe funkcje:

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);

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:

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

    adres = 2* stride * thid;
    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):

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:

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!
Retrieved from http://home.agh.edu.pl/~socha/pmwiki/pmwiki.php/CUDA/CUDA
Data ostatniej modyfikacji: 11.06.2011, 17:37