Část 3 CUDA VS. OPENCL, OPENACC

Část 3

CUDA VS. OPENCL, OPENACC

Jazyky a API pro výpočty na GPU Jazyky pro programování shaderů • HLSL – DirectX • GLSL – OpenGL • Cg – překlad do OpenGL i DirectX Nadstavby – rozšíření o proudové zpracování dat • BrookGPU / Brook+ • RapidMind • PeakStream Jazyky pro obecné výpočty na GPU (GPGPU) • CUDA • OpenCL • MS DirectCompute • OpenACC 2

CUDA (Compute Unified Device Architecture) - 2007 CUDA = API + runtime prostředí + podpora v HW • Aplikační rozhraní (API) – CUDA C – rozšíření jazyka C – CUDA driver (low-level)

•

Runtime prostředí – spouštění kernelů – přesuny dat

•

Přímá podpora v hardwaru – Architektura SIMT (Single-Instruction, Multiple-Thread) – Streaming Multiprocessor se sdílenou pamětí

3

OpenCL (Open Computing Language) 2008 Standard od Khronos Group pro paralelní výpočty • nezávislost na konkrétní hardwarové platformě • Založen na C99 obohacený o podporu pro paralelismus (datový i programový) • Podpora heterogenních systémů (tj. multi-core CPU + GPU, další typy procesorů – Xeon Phi, Cell, DSP) OpenCL framework se skládá ze 2 částí: • OpenCL C – rozšíření jazyka C • OpenCL runtime API Implementace závislá na konkrétním výrobci grafických karet a procesorů: • Překladač • Runtime prostředí

4

MS Direct Compute • Microsoft DirectCompute je API pro podporu GPGPU. • Podporované OS jsou Microsoft Windows Vista a Windows 7 a novější. • DirectCompute je součástí Microsoft DirectX API (verze 10 a 11). • Rozhraní jsou shodná nebo obdobná jako – OpenCL, – OpenGL, – CUDA.

5

OpenACC • OpenACC (Open Accelerator): nový (představen 2012) standard pro práci s akcelerátory: – – – –

GPU Nvidia GPU AMD/ATI Xeon Phi atd.

• Přístup podobný OpenMP => vyšší úroveň • Optimalizace je přenechána kompilátoru => efektivita ? • Větší podpora jen v komerčních produktech: – Portland Group (PGI), Accelerator Compiler – CAPS , HMPP Workbench – Cray Corporation, Compilation Environment

• Nekomerční podpora: – Podpora v rámci LLVM – Pracuje se na podpoře v rámci GNU GCC

6

C++ AMP • (2012?) C++ Accelerated Massive Parallelism (C++ AMP) je programovací model pro datově paralelní úlohy • C++ AMP je knihovna implementována na základě DirectX 11 a otevřené specifikace od fy Microsoft • HCC kompilátor kompiluje kód do: – OpenCL, – Standard Portable Intermediate Representation (SPIR), – HSA Intermediate Language (HSAIL)

CUDA (Compute Unified Device Architecture) CUDA je architektura pro provádění paralelních výpočtů, která definuje: • Programming model – vlákna (threads), blok (block) a mřížka (grid) • Memory model – registry, lokální, sdílená a globální paměť • Execution model – spouštění vláken a jejich mapování na HW Vývojové prostředí zahrnuje: • Driver • Runtime – spouštění C funkcí na GPU • Toolkit – překladač, debugger, profiler • Knihovny – CUFFT, CUBLAS, … • SDK – dokumentace + ukázky kódu CUDA je podporována na všech grafických procesorech NVIDIA (CUDA enabled-GPUs) počínaje čipem G80 (Geforce 8800) CUDA programy je možno psát v jazyku C/C++ nebo Fortran.

CUDA základní pojmy • CPU (označované jako host – hostitel) • Kernel = část kódu, kterou chceme provádět paralelně na GPU (def. jako funkce) • Vlákno (thread) = instance kernelu • GPU (označované jako device – zařízení) je tvořeno multiprocesory. • Streaming Multiprocesor (dále jen SM) se skládá z několika (např. 8 u G80) procesorů (jader = SP). • SM provádějí bloky vláken. • Warp je skupina vláken spouštěná najednou, např. 32 na GT200 • GPU má SIMT architekturu = všechna vlákna v rámci warpu jsou ovládány jednou instrukční jednotkou = provádí stejnou instrukci (kromě podmíněných příkazů) • Vlákna během výpočtu přistupují k různým druhům pamětí

CUDA capabilities – Compute capability 1.0 (Tesla) • G80

– Compute capability 1.1 • G86, G84, G98, G96, G96b, G94, G94b, G92, G92b

– Compute capability 1.2 • GT218, GT216, GT215

– Compute capability 1.3 • GT200, GT200b

– Compute capability 2.0 (Fermi) • GF100, GF110

– Compute capability 2.1 • GF108, GF106, GF104, GF114, GF116

– Compute capability 3.0 (Keppler) • GK104, GK106, GK107

– Compute capability 3.5 • GK110

– Compute capability 5.0 (Maxwell) • GM107

– Výpis CUDA-enabled grafických karet a jejich CC možno nalézt na: http://developer.nvidia.com/cuda-gpus 10

CUDA limity Platí stále: • Velikost warpu je 32. • Velikost konstantní paměti je 64 KB. • Cache pro konstantní paměť na SM je 8 KB. • Cache pro textury na SM je mezi 6 KB a 8 KB.

11

Compute capability 1.0 • • • • • • • •

Maximální velikost sdílené paměti na jeden SM je 16 KB. Maximální počet vláken na blok je 512. Maximální počet rezidentních bloků na SM je 8. Maximální počet rezidentní warpů na SM je 24. Počet 32bitových registrů na SM je 8K (8*1024). Počet bank sdílené paměti je 16. Velikost lokální paměti na vlákno je 16 KB. Maximální počet instrukcí pro kernel = 2 millióny

Compute capability 1.1 • Atomické instrukce 32bit int v globální paměti 12

Compute capability 1.2 • • • • •

Atomické instrukce 64bit int v globální paměti Atomické instrukce 32bit int ve sdílené paměti Podpora volby v rámci warpu Maximální počet rezidentní warpů na SM je 32. 32bitových registrů na SM je 16 K.

Compute capability 1.3 • Podpora typu double

13

Compute capability 2.0 • • • •

• • • • • • • •

Fermi architektura Atomické přičtení float v globální paměti nebo ve sdílené paměti Podpora threadfence (paměťová synchronizace). Přístupy do globální paměti používají L2 (768KB), případně L1 cache (ale možno změnit 16/48KB). Maximální počet vláken na blok je 1024. Maximální počet rezidentní warpů na SM je 48. Počet 32bitových registrů na SM je 32K (32*1024). Počet bank sdílené paměti je 32. Velikost lokální paměti na vlákno je 512 KB. Maximální velikost sdílené paměti na jeden SM je možno zvětšit na 48 KB. Zvětšena texturovací cache na SM Maximální počet instrukcí pro kernel je 512 milliónů 14

Compute capability 3.0 • • • • • •

Změna architektury (SMX), Keppler architektura Podpora pro instrukce výměny v rámci warpu. Maximální počet rezidentních bloků na SM je 16. Maximální počet rezidentní warpů na SM je 64. Počet 32bitových registrů na SM je 64K (64*1024). Zmenšena L2 cache (256/512KB)

Compute capability 3.5 • Možnost dynamického paralelismu • Zvýšení výkonu pro dvojitou přesnost • Zvětšena L2 cache (1536KB) 15

Compute capability 5.0 • • • •

Maxwell architektura Maximální počet rezidentních bloků na SM je 32. Změna architektury (SMM). Maximální velikost sdílené paměti na jeden SM je možno zvětšit na 64 KB. • Zvětšena texturovací cache na SM • Zvětšena L2 cache (2048KB)

16

Maximální výkonnost SM(X,M) Compute Capability

1.0 1.1 1.2

1.3

2.0

2.1

3.0

3.5,3.7

5.x

float add, mul, FMA

8

8

32

48

192

192

128

double add, mul, FMA

1

1

4/16

4

8

8/64

1

32-bit integer add

10

10

32

48

160

160

128

32-bit integer mul

?

?

16

16

32

32

multiple

32-bit integer compare

10

10

32

48

160

160

64

Maximální výkonnost SM(X,M) Compute Capability 1.0 1.1 1.2

1.3

2.0

2.1

3.0

3.5,3.7

5.0

32-bit integer shift

8

8

16

16

32

32/64

64

Logical operations

8

8

32

48

160

160

128

population count

?

?

16

16

32

32

32

24-bit int multiply

8

8

Multiple

Multiple

Multiple

32-bit FP reciprocal, …

2

2

4

8

32

Multiple Multiple 32

32

Maximální výkonnost SM(X,M) Compute Capability

Type conv. from 8bit and 16-bit int to 32-bit Type conv. from and to 64-bit All other type conversions

1.0 1.1 1.2

1.3

2.0

2.1

3.0

3.5,3.7

5.0

8

8

16

16

128

128

32

Multiple

1

4/16

4

8

8/32

4

8

8

16

16

32

32

32

Budoucnost • 2. Generace Maxwell – Nepodstatné změny z hlediska CUDA

• Architektura Pascal (Volta) – 3D memory. – Unified memory – spojení adresových prostorů CPU a GPU – NVLink – nova high-speed bus mezi CPU a GPU, nahrazující PCIe; udávaná propostnost 80 and 200 GB/s.

20

CUDA 1. příklad // definice kernelu __global__ void HelloWorld() { //proveden kernelem = na GPU printf(“Hello World!\n”); } int main() { // volani kernelu z funkce main HelloWorld <<<1, 1>>>(); }

Musí být void

Počet vláken v bloku Počet bloků

CUDA 2. příklad __global__ void HelloWorld2() { //identifikační číslo vlákna int i = threadIdx.x; printf(“Hello World from thread %i !\n”,i); } int main() { // volání kernelu z funkce main HelloWorld2 <<<1, 10>>>(); }

Číslo v rámci bloku

Počet vláken v bloku Počet bloků

CUDA 3. příklad __global__ void PrintInt(int *vstup) { int i = threadIdx.x; printf(“Value at %i = %i\n”,i, vstup[i]); } int main() { int i, a[N]; for(i=0;i>>(a); }

!!! Ukazatel předán, ale chyba: kernelu je předána adresa v rámci CPU nikoliv v rámci GPU !!!

Počet vláken v bloku Počet bloků

CUDA 4. příklad I __global__ void VecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { N=A.size(); VecAdd<<<1, N>>>(A, B, C); }

Číslo v rámci bloku (každé vlákno sečte unikátní položky)

Počet vláken v bloku (omezeno !!!) Počet bloků

CUDA 4. příklad II int main( void ) { int N,i,j; float *hostC, *hostA,*hostB; float *devC,*devA,*devB;

Deklarace ukazatelů V rámci CPU i GPU

Alokace paměti na GPU

cudaMalloc( (void**)&devA, N * sizeof(float) ) ; cudaMalloc( (void**)&devB, N * sizeof(float) ) ; cudaMalloc( (void**)&devC, N * sizeof(float) ) ; cudaHostAlloc( (void**)&hostA, N * sizeof(float), cudaHostAllocDefault ) ; cudaHostAlloc( (void**)&hostB, N * sizeof(float), cudaHostAllocDefault ) ; cudaHostAlloc( (void**)&hostC, N * sizeof(float), cudaHostAllocDefault ) ; Alokace paměti na CPU, Možno I pomocí malloc nebo new

CUDA 4. příklad III Init(hostA); Init(hostB);

Kopírování CPU -> GPU

cudaMemcpy(devA,hostA,sizeof(float)*N,cudaMemcpyHostToDevice); cudaMemcpy(devB,hostB,sizeof(float)*N,cudaMemcpyHostToDevice); VecAdd<<<1,N>>>( devA, devB, devC); cudaDeviceSynchronize();

Volání kernelu Čekání na dokončení kernelu

cudaMemcpy(hostC,devC,sizeof(float)*N,cudaMemcpyDeviceToHost);

Kopírování GPU -> CPU

CUDA 4. příklad IV cudaFreeHost( hostA ) ; cudaFreeHost( hostB ) ; cudaFreeHost( hostC ) ; cudaFree( devA ) ; cudaFree( devB ) ; cudaFree( devC ) ; }

Uvolnění paměťových bloků na CPU

Uvolnění paměťových bloků na GPU

Práce s globální pamětí cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, enum cudaMemcpyKind kind ); • kopírování dat mezi GPU↔CPU (cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice) cudaError_t cudaMalloc ( void **devPtr, size_t size ); • alokace paměti na GPU pro uložení dat – data nejsou přímo přístupná na CPU→kopírování dat pomocí cudaMemcpy() cudaError_t cudaFree ( void *devPtr ); • GPU nemůže přistupovat do paměti CPU a opačně • uvolnění alokované paměti na straně GPU

Zjištění dostupných zařízení cudaError_t cudaGetDeviceCount ( int *count ); • vrátí počet zařízení (tj. CUDA enabled GPU), pokud žádné neexistuje, pak vrátí cudaErrorNoDevice cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device ); vrátí informace (struktura cudaDeviceProp) o zařízení s číslem device: • jméno zařízení, • Počet SM • velikost globální paměti, • max. počet vláken na blok, apod.

Nvidia GPU Device

• Složeno z několika „Streaming multiprocessors“ (SMs) • Dále části společné: – L2 cache – Rozhraní k paměti – apod.

Multiprocessor N

Multiprocessor 2 Multiprocessor 1 Shared Memory Registers

Processor 1

Registers

Processor 2

Registers

…

Instruction Unit Processor M

Constant Cache Texture Cache

Device memory

30

Streaming multiprocessor Streaming multiprocessor (SM) obsahuje: • Instrukční cache • Jednotky pro obsluhu instrukcí (fetchdispatch) • Sdílená paměť • Registrové pole • Registry pro uložení kontextu vláken • Daný počet SP • Daný počet SFU (special function unit)

Instruction Cache Warp Scheduler

Warp Scheduler

Dispatch Unit

Dispatch Unit Register File 32768 x 32bit

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST

SFU

SFU

SFU

SFU

Interconnect Memory L1 Cache / 64kB Shared Memory L2 Cache

31

CUDA core Dispatch Port

Operand Collector

FP Unit

Int Unit

Result Queue 32

2D architektura = škálovatelnost • Každý GPU se může lišit – Frekvencí GPU – Pamětí a velikostí GPU paměti – Počtem SM – Počtem jader na SM – Schopnostmi (CUDA capabilities) Protože jednotlivé SM pracují víceméně nezávisle, je snadné zvyšovat výkon GPU zvýšením počtu SM GPU využívá masivní paralelismus na několika úrovních: • V rámci GPU existuje několika procesorů (SM=Streaming Multiprocessor), které jsou pouze volně vázané. • V rámci SM je prováděna paralelně jedna instrukce pro danou skupinu vláken (tzv. warp=32 vláken), SIMT přístup. • V rámci SM se v provádění střídají jednotlivé skupiny vláken (prokládání warpů).

Provádění GPU kódu I Vlákna jsou sdružována do bloků a bloky do mřížky (2D nebo 3D) Vlákna v rámci bloku: • lze synchronizovat (pomocí bariéry) • mohou sdílet data pomocí sdílené paměti Programátor zadá (téměř libovolně) počet bloků a počet vláken v bloku (2D z hlediska programátora) To jak přesně bude kód prováděn udává execution model. Ten se stará o korektní a efektivní mapování na dané HW. V HW totiž existuje daný počet SM a daný počet SP na jeden SM. (2D z hlediska HW)

Provádění GPU kódu II • Každý blok vláken je prováděn na právě jednom SM. Na každý SM může být namapováno více bloků najednou (až 8 u GT200) • Blok vláken je rozdělen do několika warpů (podle čísel vláken). • Centrální distribuce bloků metodou round-robin. • V každém časovém kroku plánovač (warp scheduler) každého SM (nezávisle na ostatních SM) naplánuje k provedení připravený warp (který má všechna data a neprovádí bariéru) z určité skupiny tzv. rezidentních warpů.

Provádění GPU kódu 1.

Před vlastním provedením jedné instrukce dojde k přepnutí kontextu: pro všechna vlákna aktuálního warpu se provede nahrání obsahů patřičných registrů z registrového pole SM. 2. Vlákna warpu pak provedou paralelně jednu instrukci protože jde o SIMT architekturu. 3. Přepnutí kontextu II: pro všechna vlákna aktuálního warpu se provede uložení obsahů patřičných registrů do registrového pole SM. • Toto plánování je akcelerováno pomocí hardware GPU proto má nulovou režii. GPU ale vždy vybírá z určité skupiny rezidentních warpů, jejíž velikost je hardwarově omezena. – – – –

CC 1.0: maximálně 24 rezidentních warpů CC 1.2: maximálně 32 rezidentních warpů CC 2.0+: maximálně 48 rezidentních warpů …

Provádění GPU kódu II • CPU(1 core): 1. Ins. 0 Thr. 0 2. Ins. 1 Thr. 0 3. Ins. 2 Thr. 0 4. Ins. 3 Thr. 0 5. …. 10000. Ins. 9999 Thr. 0 10001. Ins. 0 Thr. 1 10002. Ins. 1 Thr. 1 10003. Ins. 2 Thr. 1

Přepnutí kontextu

• GPU(1 SM): 1. Ins. 0 Warp 0 2. Ins. 0 Warp 1 3. Ins. 0 Warp 2 4. Ins. 0 Warp 3 5. …. 32. Ins. 1 Warp 0 33. Ins. 1 Warp 1 34. Ins. 1 Warp 2 35. Ins. 1 Warp 3

Provádění GPU kódu III • Výhody – Odstranění/ zmenšení vlivu latencí (cache misses, stalls) – Není nutné implementovat Out-of-Order execution a jiné mechanismy

• Nevýhody: – – – –

Nižší takt Velké registrové pole (+omezení reg. na vlákno) Velký počet pomocných registrů Komplikovanější přepínání kontextu vláken (registry/zásobník apod.)

Ukončení • • • •

Pokud je kód pro dané vlákno ukončen, vlákno přestane vykonávat činnost. Pokud jsou ukončeny všechna vlákna ve warpu, je warp ukončen. Pokud jsou ukončeny všechny warpy v bloku, je blok ukončen. Pokud jsou ukončeny všechny bloky v kernelu, je kernel ukončen.

39

Modifikátory funkcí 1.

__global__ : 1. voláno z CPU kódu, 2. nemožno z GPU kódu 3. Návratový typ void

2.

__device__ : 1. Voláno z ostatních GPU funkcí, 2. Nemožno z CPU kódu

3.

__host__ : – Prováděno na CPU, – Voláno z CPU

• __host__ a __device__ mohou být kombinovány – Kompilátor pak vytvoří jak CPU tak GPU kód

Bloky kontra vlákna I Je dobré mít velký počet vláken v bloku • Synchronizace a předávání dat pomocí sdílené paměti je možno jen v rámci jednoho bloku! • Díky přeplánování mohou být amortizovány velké latence. Ale: • Blok může obsahovat maximálně – 512 threadů pro CC 1.X – 1024 threadů pro CC 2.0

•

Každý SM může (současně) provádět 8 bloků (16 pro CC 3.0, 32 pro CC 5.0), ale skutečný počet závisí na paměťových požadavcích na registry a sdílenou paměť = např. všechny bloky na SM mohou alokovat – – – – –

•

max. 16 kB sdílené paměti (pro CC 1.X), max. 48 kB sdílené paměti (pro CC 2.0), max. 112 kB sdílené paměti (pro CC 3.7) max. 64 kB sdílené paměti (pro CC 5.0), max. 96 kB sdílené paměti (pro CC 5.2).

Stejně tak je limitován počet registrů pro vlákno.

GPGPU II: Programování v CUDA

Bloky kontra vlákna II Omezení bloků: • Bloky by měly být nezávislé • Být navrženy tak, aby každé pořadí jejich vyhodnocení bylo korektní • Mohou běžet paralelně nebo sekvenčně • Data mohou být sdílena mezi bloky ale problém se synchronizací

Ošetření chyb 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 ); Makro pro zpracování chyb }} #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

…. HANDLE_ERROR( cudaMalloc( (void**)&devA, N * sizeof(float) ) ); ….

Výpis případné chyby

Zabudované vektorové typy Mohou být použity v GPU nebo CPU kódu – [u]char[1..4], – [u]short[1..4], – [u]int[1..4], – [u]long[1..4], – float[1..4] •

Struktury přístupné nikoliv pomocí indexů ale pomocí položek x, y, z, w:

Např. uint4 param; int y = param.y; • Speciální typ dim3, založen na typu uint3 – používán ke specifikaci dimenzí – Defaultní hodnota (1,1,1)

Mřížka bloků • Bloky mi mohou obecně tvořit maximálně 2D (pro CC 1.X) nebo 3D (pro CC 2.0) mřízku = grid • Obdobně vlákna v rámci bloku mi mohou obecně tvořit (maximálně 3D) matici. • Ale mřížka i matice jsou převedeny do 1D, takže nemají velký význam kromě vyššího komfortu pro programátora při přístupu do vícedimenzionálního pole.

Mapování

Automaticky definované proměnné I • Všechny __global__ a __device__ funkce mají přístup k těmto automaticky definovaným proměnným • dim3 gridDim; – dimenze mřížky v blocích

• dim3 blockDim; – Dimenze bloku ve vláknech

• dim3 blockIdx; – Index bloku v gridu

• dim3 threadIdx; – index vlákna v bloku

• Častý přepočet: int idx = blockDim.x * blockIdx.x + threadIdx.x;

Automaticky definované proměnné II • Obsahují informace o konfiguraci mřížky a také jednoznačné identifikátory jednotlivých vláken v rámci bloku a bloků v rámci mřížky: • gridDim udává velikost mřížky, tj. počet bloků v jednotlivých dimenzích. Maximální velikost kterékoliv dimenze je 2^16-1 (pro CC 1.0--2.1) nebo 2^31-1 (pro CC 3.0). • blockDim určuje velikost bloku v jednotlivých dimenzích. Blok je chápán jako třírozměrné pole vláken, jejichž maximální počet v rámci bloku je omezen na 512 pro CC 1.X a 1024 pro CC 2.0, max. velikost z-dimenze je 64).

Automaticky definované proměnné III • Každé vlákno je jednoznačně identifikováno pomocí: – čísla bloku v rámci mřížky blockIdx – čísla vlákna v rámci bloku threadIdx

• warpSize definuje počet vláken tvořících warp (velikost warpu je zatím rovna 32).

Mapování I __global__ void MujKernel(int i) {// v tomto kernelu bude mit: // gridDim.x stále hodnotu 5 // gridDim.y stále hodnotu 7 // gridDim.z stále hodnotu 1

// blockDim.x stále hodnotu 4 // blockDim.y stále hodnotu 3 // blockDim.z stále hodnotu 2

int parametr; dim3 gridRes(5,7,1); ); //gridRes mi určuje třírozměrnou mřížku bloků dim3 BlockRes(4,3,2); //blockRes mi určuje třírozměrné pole vláken MujKernel<<>>(parametr);

50

Mapování II __global__ void MujKernel(int i) {// v tomto kernelu bude mit: // blockIdx.x hodnotu mezi <0,gridDim.x) tj. mezi 0 a 4 // blockIdx.y hodnotu mezi <0,gridDim.y) tj. mezi 0 a 6 // blockIdx.z hodnotu mezi <0,gridDim.z) tj. 0 // threadIdx.x hodnotu mezi <0,blockDim.x) tj. mezi 0 a 3 // threadIdx.y hodnotu mezi <0,blockDim.y) tj. mezi 0 a 2 // threadIdx.z hodnotu mezi <0,blockDim.z) tj. mezi 0 a 1 }

int parametr; dim3 gridRes(5,7,1); //gridRes mi určuje třírozměrnou mřížku bloků dim3 BlockRes(4,3,2); //blockRes mi určuje třírozměrné pole vláken MujKernel<<>>(parametr); 51

Mapování III // převod 2D indexu vlákna na linearní index (v rámci bloku ) int idx = blockDim .x* threadIdx .y + threadIdx .x; // převod 3D indexu vlákna na lineární (v rámci bloku ) int idx = blockDim .x*( blockDim .y* threadIdx .z + threadIdx .y) + threadIdx .x; // výpočet globálního indexu vlákna v rámci 2D mřížky int column = blockDim .x * blockIdx .x + threadIdx .x; int row = blockDim .y * blockIdx .y + threadIdx .y;

Součet vektorů I __global__ void VecAdd2(float* A, float* B, float* C,int N) { int i; for(i=0;i>>(A, B, C, N); } Sekvenční !!!!

Součet vektorů II __global__ void VecAdd(float* A, float* B, float* C) { int i = threadIdx.x; C[i] = A[i] + B[i]; } int main() { VecAdd2<<<1, N>>>(A, B, C); }

Číslo v rámci bloku (každé vlákno sečte unikátní položky)

Počet vláken v bloku (omezeno N≤1024) Počet bloků

Součet vektorů III Přepočet na unikátní číslo v rámci mřížky (každé vlákno sečte unikátní položky)

__global__ void VecAdd3(float* A, float* B, float* C, int N) { int i = blockDim .x * blockIdx .x + threadIdx .x; if (i>>(A, B, C); } Celkový počet vláken ≥ N

Násobení matic __global__ void Mul_IJ( float *inA, float *inB, float *outC, int N ) { int i=blockIdx.x; int j=threadIdx.x; Index bloku int k; Index v rámci bloku float s=0.0; for(k=0;k
Provedení skalárního součinu

Mul_IJ<<>>( devA, devB, devC, N ); Volání kernelu

Trojúhelníková matice __global__ void TrojMatice( float *A, int N ) { int i=blockIdx.x; int j=threadIdx.x; if (j>i) return;

Index bloku Index v rámci bloku

A[i*n+j]*=2.0; }

Elegantní, ale trochu neefektivní

TrojMatice<<>>( devA, N ); Volání kernelu pro zpracování troj. matice

Synchronizace s hostitelem • Spuštení kernelů je asynchronní operace – Vrací řízení CPU ihned jak je to možné – Pokud je třeba počkat na dokončení práce všech CUDA volání: cudaThreadSynchronize() nebo spíše cudaDeviceSynchronize() – kernel je spuštěn po dokončení všech předchozích CUDA voláních

• Např. operace kopírování cudaMemcpy() je synchronní, je spuštěna po dokončení všech předchozích CUDA volání

Synchronní CUDA volání • • • •

cudaMalloc(void **pointer, size_t nbytes) cudaMemset(void *pointer, int value, size_t count) cudaFree(void *pointer) cudaMemcpy(void *dst, void *src, size_t nbytes, enum cudaMemcpyKind direction);

Asynchronní CUDA volání Vzhledem k hostiteli: • Spuštění kernelu v default streamu • cudaMemcpy*Async • cudaMemset*Async • cudaMemcpy v rámci stejného device • HostToDevice cudaMemcpy pro 64kB nebo menší blok Vzhledem k device: • Pomocí více streamů (bude uvedeno později)

Synchronizace v rámci bloku I • Většina vláken běží asynchronně pouze vlákna v rámci jednoho warpu musí běžet plně synchronně. • Synchronizace (bariéra) v rámci celého bloku pomocí built-in funkce __syncthreads() • Žádné vlákno nemůže překročit bariéru, pokud ji nedosáhnou všechny vlákna bloku – Časté použití: synchronizace přístupu ke sdílené paměti – Předcházení RAW, WAR, WAW hazardů

Synchronizace v rámci bloku II • Je třeba velké opatrnosti při použití např. v podmínce //dimBlock.x = 256 if (threadIdx.x < 128) { __syncthreads(); // CHYBA ! } • Možno jen v podmínce, která je platná pro všechny vlákna v bloku: if (blockIdx.x < 128) { ... } !!! Jinak hrozí deadlock nebo jiné nepředvídatelné chování !!!

CUDA typy pamětí I • Registry – – – –

Jen pro jedno vlákno Umístěna na čipu, velmi rychlá Počet použitých registrů pro každé vlákno je dán kernelem Počet použitých registrů omezuje počet rezidentních warpů a bloků na jednom SM (všechny se musí vejít do reg. pole daného SM)

CUDA typy pamětí II • Globální – Přístupná i pro hostitele – Přístupná pro všechny bloky – Umístěna mimo čip a pro • CC 1.X nemá cache => velká latence • CC 2.0 při přístupu využívána cache

– Pro maximální využití přenosové rychlosti je nutný sdružený přístup (coalesced access) – Paralelní přístupy serializovány – Pro implementaci není nutný žádný koherenční protokol

CUDA typy pamětí III • Sdílená (shared) – Přístupná pro všechny vlákna ale jen uvnitř jednoho bloku – Umístěna na čipu, velmi rychlá, neobsahuje cache – Rozdělena na: • Pro CC 1.X 16 bank (prokládaně po 32b) • Pro CC 2.0 32 bank (prokládaně po 32b)

– Současný přístup k různým adresám v jedné bance vede k serializaci – Při čtení ze stejné adresy je hodnota distribuována pomocí operace broadcast – Její velikost omezuje počet rezidentních bloků na jednom SM (všechny se musí vejít do sdílené paměti daného SM)

CUDA typy pamětí IV • Lokální (lokální) – Přístupná pro jedno vlákno – Část globální paměti (rychlost?) – Jako možné rozšíření registrového pole (tzv. register spilling)

CUDA typy pamětí V • Pro konstanty (constant) – – – – –

Přístupná pro všechny vlákna Umístěna mimo čip, ale obsahuje vlastní cache => rychlost? Část globální paměti určená jen pro čtení Pro implementaci není nutný žádný koherenční protokol (jen pro čtení) Speciální operace cudaMemcpyToSymbol pro načtení dat do konstantní paměti

CUDA typy pamětí VI • Texturovací (texture) – – – –

Přístupná pro všechny vlákna Umístěna mimo čip, ale obsahuje vlastní cache => rychlost? Část globální paměti určená jen pro čtení (speciální 2D kešování) Pro implementaci není nutný žádný koherenční protokol (jen pro čtení)

CUDA typy pamětí VII • Pamět jen pro čtení – – – –

Od CC 3.0 Přístupná pro všechny vlákna Část globální paměti určená jen pro čtení Umístěna mimo čip, ale používá texturovou cache i pro jiné datové objekty než textury – Přistupované proměnné musí být deklarována s modifikátory const a __restrict__ – Pro implementaci není nutný žádný koherenční protokol (jen pro čtení)

CUDA typy pamětí Druh paměti

Přístup

Umístění

Operace

Kešovaná

Registry

1 vlákno

Na čipu

R+W

NE

Lokální

1 vlákno

DRAM

R+W

NE/ANO

Sdílená

Všechny vlákna bloku

Na čipu

R+W

NE

Globální

Všechna vlákna a host

DRAM

R+W

NE/ANO

Texturovací

Všechna vlákna a host

DRAM

R

ANO

Pro konstanty

Všechna vlákna a host

DRAM

R

ANO

Jen pro čtení

Všechna vlákna a host

DRAM

R

ANO

Jak specifikovat umístění proměnných (GPU kód) • __device__ – Uloženo v globální paměti – Alokováno pomocí cudaMalloc – Přístupná pro všechny vlákna

• __shared__ – Ve sdílené paměti – Alokováno pomocí execution configuration nebo v době kompilace

• __constant__ – V paměti konstant

• Neoznačené proměnné: – Skaláry a zabudované vektorové typy jsou uloženy v registrech – Pole v lokální paměti (registry nejsou adresovatelné)

Použití sdílené paměti • V čase kompilace __global__ void kernel(…) { __shared__ float sData[256]; … } int main(void) { kernel<<>>(…); }

• Až při spuštění kernelu __global__ void kernel(…) { extern __shared__ float sData[]; … } int main(void) { smBytes = blockSize*sizeof(float); kernel<<>>(…); }

Konstantní paměť • • •

Proměnné s modifikátorem __constant__ Automaticky alokováno potřebné místo. Musíme pouze zajistit nakopírování příslušných dat z hlavní paměti do paměti konstant pomocí cudaMemcpyToSymbol.

Př. • __constant__ float constData[256]; // pole v konstantní paměti • float data[256]; // pole v paměti hostitele • •

// kopírování pole z hlavní paměti do konstantní paměti cudaMemcpyToSymbol(constData, data, sizeof(data));

Dynamicky alokovaná paměť I • void* malloc (size t size); • free (void* ptr); • Paměť je kernelem dynamicky alokována resp. vrácena zpět z haldy pevné velikosti vytvořené v globální paměti • adresa paměťového bloku velkého nejméně size bytů je zarovnána na 16 bytů • v případě neúspěchu vrací NULL • Může být používána či dealokována i jinými vláknem, než která jí původně vytvořilo.

Dynamicky alokovaná paměť II • Standardní velikost haldy je 8 MB, ale jsou k dispozici funkce, které umožňují zjistit její aktuální velikost a tuto velikost změnit: • void* cudaDeviceGetLimit (size t *size, cudaLimitMallocHeapSize); • cudaDeviceSetLimit (cudaLimitMallocHeapSize, size t size); • Možno měnit i: • cudaLimitStackSize: max. velikost zásobníku pro GPU vlákno • cudaLimitPrintfFifoSize: velikost FIFO pro printf() a fprintf().

Násobení matic (CUDA v1) I 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 ); Zpracování chyb }} #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

int main( void ) { int N,i,j;

float *hostC, *hostA,*hostB; float *devC,*devA,*devB;

Ukazatele

Násobení matic (CUDA v1) II __global__ void Mul_IJ( float *inA, float *inB, float *outC, int N ) { int i=blockIdx.x; int j=threadIdx.x; Index bloku int k; Index v rámci bloku float s=0.0; for(k=0;k
Násobení matic (CUDA v1) III 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 ); }} Makro pro zpracování chyb #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ )) int main( void ) { cudaDeviceProp prop; int whichDevice,N,i,j; cudaEvent_t start, stop; float elapsedTime; float *hostA,*hostB,*hostC; float *devA,*devB,*devC;

Vlastnosti GPU

CUDA události

Ukazatele

Násobení matic (CUDA v1) IV HANDLE_ERROR( cudaGetDevice( &whichDevice ) ); HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );

Inicializace cudaEventCreate( &start ) ; cudaEventCreate( &stop ) ;

Alokace paměti na GPU

cudaMalloc( (void**)&devA, N * N* sizeof(float) ) ; cudaMalloc( (void**)&devB, N * N *sizeof(float) ) ; cudaMalloc( (void**)&devC, N * N * sizeof(float) ) ;

Alokace paměti na CPU cudaHostAlloc( (void**)&hostA, N * N* sizeof(float), cudaHostAllocDefault ) ; cudaHostAlloc( (void**)&hostB, N * N *sizeof(float), cudaHostAllocDefault ) ; cudaHostAlloc( (void**)&hostC, N * N *sizeof(float), cudaHostAllocDefault ) ;

Násobení matic (CUDA v1) V cudaMemcpy(devA,hostA,sizeof(float)*N*N,cudaMemcpyHostToDevice); cudaMemcpy(devB,hostB,sizeof(float)*N*N,cudaMemcpyHostToDevice); Kopírování CPU -> GPU cudaEventRecord( start, 0 ) ; Volání kernelu Mul_IJ<<>>( devA, devB, devC, N ); cudaThreadSynchronize(); cudaEventRecord( stop, 0 ) ; Zjištění času

cudaEventSynchronize( stop ) ; cudaEventElapsedTime( &elapsedTime, start, stop )); printf( „GPU time taken: %g ms\n", elapsedTime ); cudaMemcpy(hostC,devC,sizeof(float)*N*N,cudaMemcpyDeviceToHost); Kopírování GPU -> CPU

Násobení matic (CUDA v1) VI cudaEventDestroy( start ) ; cudaEventDestroy( stop ) ; cudaFreeHost( hostA ) ; cudaFreeHost( hostB ); cudaFreeHost( hostC ) ;

cudaFree( devA ) ; cudaFree( devB ) ; cudaFree( devC ) ; }

Rozdělení (Diverging) warpu I • Nastává při podmíněném skoku (důsledek SIMT přístupu) • Vlákna ve stejném warpu mohou vykonávat různé větve výpočtu. Příklad: if (threadIdx.x < 11) { branch1(); } else { branch2(); }

Rozdělení (Diverging) warpu II • Warp musí provést obě větve • Warp se rozdělí v podmínce, vlákna dělají všechna stejnou instrukci (ale ty které jsou ve „špatné“ větvi nevykonávají instrukce) • Při víceúrovňovém větvení může dojít k extrému: každé vlákno ve warpu je vlastně prováděno sekvenčně • Pokud počet aktivních vláken ve warpu klesne na nulu, je ukončeno provádění tohoto warpu

Rozdělení (Diverging) warpu III Rozdělení (diverging) warpu nastává, pokud: • if/else s různou hodnotou podmínky v rámci warpu • Vlákna provádí různý počet iterací cyklu • Pokud počet vláken v bloku není násobek 32 (velikosti warpu) Rozdělení (diverging) warpu nenastává, pokud: • Jsou aktivní všechna vlákna v rámci warpu • Jsou neaktivní všechna vlákna v rámci warpu

Rozdělení (Diverging) warpu IV Problém s efektivitou může nastat při provádění tohoto kódu for(i=0;i
Rozdělení (Diverging) warpu V Předchozí kód možno nahradit takto: for(i=0;i
Nevhodné, pokud jsou podmínky jednoduché => režie operací ulož1, ulož2, ulož3 Nutnost opakovaného průchodu datovými strukturami

GPU Atomic Integer Operations • Podpora atomických operací pro typ integer – v globální paměti CUDA capabilities >= 1.1 – ve sdílené paměti CUDA capabilities >= 1.2

• Týká se těchto operací – – – –

atomicMin(), atomicMax(), atomicAdd(), atomicSub(), atomicInc(), atomicDec(), exchange ( atomicExch() ), compare and swap ( atomicCAS() ) atomicAnd(), atomicOr(), atomicXor()

• Např. __shared__ totalSum; atomicAdd(&totalSum, 1);

OpenCL (Open Computing Language) - 2008 Standard od Khronos Group pro paralelní výpočty • nezávislost na konkrétní hardwarové platformě • Založen na C99 obohacený o podporu pro paralelismus (datový i programový) • Podpora heterogenních systémů (tj. multi-core CPU + GPU, další typy procesorů – Cell, DSP) OpenCL framework se skládá ze 2 částí: • OpenCL C – rozšíření jazyka C • OpenCL runtime API Implementace závislá na konkrétním výrobci grafických karet a procesorů: • Překladač • Runtime prostředí

Převzato z [1]

Programovací jazyk OpenCL C Založen na C99, rozšíření: • Vektorové datové typy. • Datové typy a funkce podporující práci s obrázky a jejích filtrování. • Kvalifikátory adresního prostoru. • Kvalifikátory přístupových práv. • Kernelové funkce. • Přesnost čísel v plovoucí desetinné čárce dle standardu IEEE 754.

Programovací jazyk OpenCL C Omezení: • Ukazatele na funkce, pole proměnné délky a bitová pole jsou zakázaná. • Mnoho hlavičkových souborů standardní knihovny jazyka C je nedostupná. • Rekurzivní funkce nejsou povolené. • Kernelové funkce nesmějí deklarovat argumenty typu ukazatel na ukazatel ani nic vracet. • Zápisy na pole číselných typů menších než 32 bitů jsou zakázané.

Novinky I Profil „embedded“ • profil pro mobilní či vestavěná zařízení, která jsou schopná podporovat modely architektury OpenCL, ale nedisponují dostatečným výkonem pro zajištění plného rozsahu funkčnosti. • některé části standardu nepovinné (podpora 3D obrazu) • Nebo odstraněné úplně (striktní konformita s IEEE754, 64bitové číselné typy).

OpenCL 2.0 Updates and additions: • Shared virtual memory • Nested parallelism • Generic address space • Images • C11 atomics • Pipes • Android installable client driver extension

OpenCL 2.1 November 16, 2015. • the OpenCL C kernel language is replaced with OpenCL C++, a subset of C++14 • Copying of kernel objects and states • Low-latency device timer queries • Ingestion of SPIR-V code by runtime • Execution priority hints for queues • Zero-sized dispatches from host

Firmy na OpenCL • AMD/ATI, • IBM, • Intel • nVidia • Apple Nový standard 2.1 je podporován: • AMD, ARM, Intel, HPC, YetiWare

Srovnání s CUDA • OpenCL má obecnější model, díky tomu je „ukecanější“ • Je těžké napsat kód, který bude efektivní na všech OpenCL platformách • OpenCL kvůli obecnosti nedokáže plně využít všechny HW features architektury • OpenCL má vyšší režii, kód kernelu je kompilován až za běhu programu

Základní pojmy I • • • • •

Host (CUDA: host) = CPU. Device (CUDA: device) výpočetní zařízení CPU nebo GPU. Platform: systém (host + devices) spravovaný pomocí OpenCL Context: definuje celé prostředí OpenCL Kernel (CUDA: kernel) fce volaná z hosta, vykonávaná na device. Kompilované až při spuštění • Program: množina kernelů a dalších fcí

Základní pojmy II • Work-item (CUDA: thread), každý své ID. • Work-group (CUDA: thread block) skupina work-items které mohou kooperovat a komunikovat, každý své ID. Je to Ndimensional grid of work-groups, N = 1, 2 or 3). • ND-Range: popisuje velikosti dimenzí work-groups (jako Ndimensional grid of work-groups, N = 1, 2 or 3). • Compute Unit (CUDA: Streaming Multiprocessor) • Processing Element (CUDA: Streaming processor, core)

Identifikace WI • Jednotlivé WI jsou jednoznačně dány: – global id (unikátní v rámci index space) – work-group ID a local ID v rámci work-group

Identifikace • Pro identifikaci vláken v každé dimenzi: – get_global_id(dim) – get_global_size(dim)

• Nebo zjištění work-group ID and ID v rámci WG – – – –

get_group_id(dim) get_num_groups(dim) get_local_id(dim) get_local_size(dim)

get_global_id(0) = column, get_global_id(1) = row get_num_groups(0) * get_local_size(0) == get_global_size(0)

OpenCL Je definováno několik modelů • Platform Model • Execution Model • Memory Model • Programming Model Jejich účel je obdobný jako v CUDA, navíc je „platform“, který CUDA nepotřebuje

Platform Model • Každá OpenCL implementace definuje platformu, která umožnuje hostisteli využívat OpenCL zařízení

• OpenCL používá “Installable Client Driver” model – Umožňuje více platforem (od každého výrobce jednu) na jednom systému, ale mohou zde být různá omezení souběžného provozu více platforem

Platform Model • Hostitel je připojen na několik OpenCL zařízení (device) • Jedno zřízení je rozděleno na několik compute units (CU) • Každé CU je rozdělenona několik processing elements (PE) – Každý PE má vlastní program counter

Memory Model Každá pracovní jednotka (work item) má přístup do následujících pamětí: • Global Memory (CUDA: globální paměť) – hlavní paměť, která je přístupná všem pracovním jednotkám jak pro čtení, tak pro zápis • Constant Memory (CUDA: paměť pro konstanty) – oblast globální paměti jejíž obsah zůstává během spuštění kernelu konstantní • Local Memory (CUDA: sdílená paměť)– paměť sdílená všemi pracovními jednotkami v rámci pracovní skupiny (work-groups) • Private Memory (CUDA: lokální paměť a registry) – privátní paměť přístupná pouze jednotlivým pracovním jednotkám (work-items) • Private to a work-item

Paměťové objekty • Buffers – Souvislé kusy paměti – Přímý přístup (arrays, pointers, structs) – Read/write

• Images – Objekty se souřadnicemi(2D nebo 3D) – Přístup pomocí read_image() and write_image() – Buď jen pro čtení nebo jen pro zápis

Programming model • Work-groups jsou prováděny na CU( computeunits) – Není zaručena komunikace/koherence mezi různými work-groups (není obsaženo v OpenCL specifikaci)

• Synchronizace – Mezi WI v rámci WG – Mezi příkazy v kontextu příkazové fronty

Program • Program (jedná se o objekt) je kolekce OpenCL kernelů – Může to být zdrojový kód v textové formě nebo překompilovaný binární kód – Může obsahovat konstantní data a pomocné funkce

• Vytváření programu vyžaduje načtení zdrojového kódu v textové formě nebo překompilovaného binárního kódu • ke zkompilování programu je nutno: • Specifikovat cílové zařízení – Program je zkompilován zvlášť pro každé zařízení

• Zahrnutí volitelných nastavení kompilátoru • Zjištění případných chyb při kompilaci

Mapování na HW • Pro AMD (Intel ?) vícejádrové CPU – všechny CPU tvoří jedno zařízení (device) – každé jádro je jedna CU a jeden PE

• pro GPU – každé GPU tvoří zvláštní zařízení (device) – jedno VLIW jádro tvoří jeden Processing element (PE) – Jeden SIMD Engine tvoří jednu compute unit (CU)

Nvidia GPU Instruction Cache Warp Scheduler

Warp Scheduler

Dispatch Unit

Dispatch Unit Register File 32768 x 32bit

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Core

Interconnect Memory L1 Cache / 64kB Shared Memory L2 Cache

LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST LDST

SFU

SFU

SFU

SFU

CUDA core • Jeden SP = jedno jádro v CUDA je tvoří pouze ALU a FPU

Dispatch Port

Operand Collector

FP Unit

Int Unit

Result Queue

AMD GPU HW Architecture •AMD 5870 – Cypress •20 SIMD engines •16 SIMD units per core •5 multiply-adds per functional unit (VLIW processing) •2.72 Teraflops Single Precision •544 Gigaflops Double Precision

SIMD Engine • A SIMD engine consists of a set of “Stream Cores” • Stream cores arranged as a five way Very Long Instruction Word (VLIW) processor – Up to five scalar operations can be issued in a VLIW instruction – Scalar operations executed on each processing element

One SIMD Engine

One Stream Core Instruction and Control Flow

• Stream cores within compute unit execute same VLIW instruction – The block of work-items that are executed together is called a wavefront. – 64 work items for 5870

General Purpose Registers

Source: AMD Accelerated Parallel Processing OpenCL Programming Guide

Architektura AMD/ATI VLIW4 nebo VLIW5 Jedno GPU je složeno z Compute Units (CU) = SIMD engines. Každá Compute Units z proudových jader (Stream Core, SC). Každé SC obsahuje: • 3 nebo 4 Procesních elementů, • T-Procesního element, • jednotku vykonávající větvení programu (Branch Execution Unit). Jednotlivé PE a T-PE (=ALU) provádějí samotný výpočet. Každý PE a T-PE umožňuje vykonávat integer operace a FP operace v jednoduché přesností. FP operace ve dvojnásobné přesnosti lze vykonávat při spojení dvou až čtyř PE. T-PE umožňuje navíc vykonávat matematické operace jako je sinus, cosinus, logaritmus, atd.

Architektura AMD/ATI VLIW4 nebo VLIW5 SC = čtyř- nebo pěti-cestný very long instruction word (VLIW) procesor, který vykonává až čtyři (VLIW4) nebo pět (VLIW5) skalárních operací současně v jedné VLIW instrukci. Blok vláken vykonávaných paralelně nazývá Wave-Front (velikost v současné generaci je 64). Zhruba odpovídá CUDA warpu.

Architektura GCN • GCN (Graphics Core Next} • AMD/ATI opustila koncept VLIW, nová technologie Southern Islands. • První zástupci jsou karty s technologií GCN, v současné době jsou to karty Radeon HD 77xx-79xx. • Každá GCN karta je složena z několika CU. Každá CU obsahuje: – – – – –

jednotka vykonávající větvení programu (Branch Execution Unit), plánovač(e), 4 vektorových (SIMD) jednotek (Vector Units=VU), skalární jednotky (Scalar Unit), registrové pole (jak pro skalární tak pro vektorové jednotky).

OpenCL kompilace • LLVM - Low Level Virtual Machine • Kernely jsou zkomiplovány do LLVM IR • Open Source Compiler – Platform, OS independent – Multiple back ends • http://llvm.org

Fáze v OpenCL 1. 2. 3. 4. 5. 6. 7. 8. 9. 10. 11. 12. 13. 14.

zjištění dostupných platforem výběr vhodného zařízení vytvoření kontextu vytvoření fronty příkazů vytvoření paměťových objektů pro uložení dat načtení kernelů ze souboru (případně definice pomocí stringu) vytvoření objektu pro uložení programu (program object) vytvoření spustitelné verze programu (překlad a linkování) vytvoření objektu kernelu (kernel object) nastavení parametrů kernelu překopírování dat do zařízení spuštění kernelu překopírování výsledků ze zařízení uvolnění alokovaných prostředků (objektů)

Fáze v CUDA

1. výběr vhodného zařízení 2. alokace paměťových objektů pro uložení dat 3. překopírování dat do zařízení 4. spuštění kernelu 5. překopírování výsledků ze zařízení 6. uvolnění alokovaných prostředků (objektů)

Příkazová fronta (Command-Queue) • Objekt, kam jsou uloženy OpenCL příkazy. Více těchto front umožnuje zařízení provádět (paralelně) více příkazů bez potřeby synchronizace. • Typy příkazů: – spuštění kernelu na výpočetních elementech zařízení – datové přenosy z/do paměťových objektů a přenosy mezi paměťovými objekty, mapovaní paměťových objektů do paměti hosta – synchronizační příkazy

• Příkazy ve frontách se mohou provádět dvěma způsoby: – In-order Execution – sériové provádění příkazů ve frontě (předchozí příkaz musí být dokončen než je spuštěn další) – Out-of-order Execution – nečeká se na dokončení předchozích příkazů, ale obvykle vyžaduje synchronizační příkazy (bariera, události)

Kód kernelu __kernel void vector_add_gpu ( __global const float* src_a, __global const float* src_b, Vstupní parametry __global float* res, const int num) { Číslo work-item const int idx = get_global_id(0); if (idx < num) res[idx] = src_a[idx] + src_b[idx]; } • Návratový typ kernelu je vždy void. • Všechny kernely musí být v souborech ".cl" files, které obsahují jen OpenCL kód.

Fáze 1) cl_int error = 0; cl_platform_id platform; error = oclGetPlatformID(&platform); if (error != CL_SUCCESS) { cout << "Error getting platform id: " << errorMessage(error) << endl; exit(error); } • Zjistí identifikátory dostupných platforem • Informace o konkrétní platformě (jméno, podporovaná verze OpenCL, výrobce, podporované extenze) mohou být zjištěny pomocí příkazu clGetPlatformInfo(…)

Fáze 2) cl_device_id device; error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); if (err != CL_SUCCESS) { cout << "Error getting device ids: " << errorMessage(error) << endl; exit(error); } • Typ zřízení může být CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_DEFAULT, CL_DEVICE_TYPE_ALL • informace o konkrétním zařízení (typ, počet výpočetních jednotek, max.velikost pracovní skupiny, …) mohou být zjištěny pomocí příkazu clGetDeviceInfo(…)

Fáze 3) cl_context context; context = clCreateContext(0, 1, &device, NULL, NULL, &error); if (error != CL_SUCCESS) { cout << "Error creating context: " << errorMessage(error) << endl; exit(error); } Kontext může zahrnovat jedno nebo více zařízení

Fáze 4) cl_command_queue queue; queue = clCreateCommandQueue(context, device, 0, &error); if (error != CL_SUCCESS) { cout << "Error creating command queue: " << errorMessage(error) << endl; exit(error); } • 3. parametr je požadované vlastnosti fronty je logický součet hodnot CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE a CL_QUEUE_PROFILING_ENABLE

Fáze 5) a 11) const int size = 1234567; float* src_a_h = new float[size]; float* src_b_h = new float[size]; float* res_h = new float[size]; for (int i = 0; i < size; i++) { src_a_h = src_b_h = (float) i; } const int mem_size = sizeof(float)*size; cl_mem src_a_d = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size, src_a_h, &error); cl_mem src_b_d = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size, src_b_h, &error); cl_mem res_d = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size, NULL, &error);

Vytvoření bufferu cl_mem clCreateBuffer (cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret) Velikost v B

V kterém kontextu

vrací objekt typu cl_mem odkazující na paměť alokovanou na zařízení položka flags obsahuje bitově zapsané tyto možnosti: • CL_MEM_READ_WRITE • CL_MEM_WRITE_ONLY • CL_MEM_READ_ONLY • CL_MEM_USE_HOST_PTR • CL_MEM_ALLOC_HOST_PTR • CL_MEM_COPY_HOST_PTR – kopíruje paměť z host_ptr

Fáze 6), 7), 8) // Uses NVIDIA helper functions to get the code string and it's size (in bytes) size_t src_size = 0; const char* path = shrFindFilePath("vector_add_gpu.cl", NULL); const char* source = oclLoadProgSource(path, "", &src_size); cl_program program = clCreateProgramWithSource(context, 1, &source, &src_size, &error); assert(error == CL_SUCCESS); // Builds the program error = clBuildProgram(program, 1, &device, NULL, NULL, NULL); assert(error == CL_SUCCESS); • kernel lze definovat v programu jako string nebo je nutno kód kernelu načíst ze souboru (buď jako zdrojový kód nebo jeho binární verzi) • spustitelná verze programu je vytvořena voláním OpenCL překladače (kompilace + slinkování)

Fáze 9) char* build_log; size_t log_size; // First call to know the proper size clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); build_log = new char[log_size+1]; // Second call to get the log clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL); build_log[log_size] = '\0'; cout << build_log << endl; delete[] build_log; // Extracting the kernel cl_kernel vector_add_kernel = clCreateKernel(program, "vector_add_gpu", &error); assert(error == CL_SUCCESS); • spustitelná verze programu může obsahovat více kernelů, musíme pro každý kernel, který chceme spouštět, vytvořit jeden objekt kernelu (obsahuje jméno spouštěné funkce a popis jejích parametrů)

Fáze 10 // Enqueuing parameters // Note that we inform the size of the cl_mem object, not the size of the memory pointed by it error = clSetKernelArg(vector_add_k, 0, sizeof(cl_mem), &src_a_d); error |= clSetKernelArg(vector_add_k, 1, sizeof(cl_mem), &src_b_d); error |= clSetKernelArg(vector_add_k, 2, sizeof(cl_mem), &res_d); error |= clSetKernelArg(vector_add_k, 3, sizeof(size_t), &size); assert(error == CL_SUCCESS); // Launching kernel • před spuštěním kernelu je nutné nastavit hodnotu každého parametru

Fáze 12), const size_t local_ws = 512; // Number of work-items per work-group // shrRoundUp returns the smallest multiple of local_ws bigger than size const size_t global_ws = shrRoundUp(local_ws, size); // Total number of work-items error = clEnqueueNDRangeKernel(queue, vector_add_k, 1, NULL, &global_ws, &local_ws, 0, NULL, NULL); assert(error == CL_SUCCESS); • •

5.parametr je celkový počet pracovních jednotek (pole obsahující počet pracovních jednotek v jednotlivých dimenzích) 6.parametr je dimenze pracovní skupiny (pole obsahující počet pracovních jednotek v jednotlivých dimenzích)

Fáze 13) // Reading back float* check = new float[size]; clEnqueueReadBuffer(queue, res_d, CL_TRUE, 0, mem_size, check, 0, NULL, NULL); •

CL_TRUE = blokující synchronní čtení (čeká se na jeho dokončení)

Fáze 14) // Cleaning up delete[] src_a_h; delete[] src_b_h; delete[] res_h; delete[] check; clReleaseKernel(vector_add_k); clReleaseCommandQueue(queue); clReleaseContext(context); clReleaseMemObject(src_a_d); clReleaseMemObject(src_b_d); clReleaseMemObject(res_d);

OpenMP 4.0 • Introduced in 2013, standardization was done by: – AMD, Cray, Fujitsu, HP, IBM, Intel, Nvidia, etc.

• Similar to, but not the same as OpenACC directives. • Support for more than just loops • Not GPU specific – suitable for Xeon Phi or DSPs, for example

• Fully integrated into the rest of OpenMP • Supported compilers: – GCC: GCC 4.9 supports OpenMP 4.0 for C/C++, GCC 4.9.1 also for Fortran. GCC 5 adds support for Offloading – LLVM 3.8

– Intel C++ Composer XE 2013

131

What is new in OpenMP 4.0 • • • • • •

Support for accelerators (or heterogeneous devices) Thread affinity support SIMD support for vectorization Thread cancellation Fortran 2003 support Extended support for – Tasking (groups, dependencies, abort) – Reductions (i.e. User Defined Reductions) – Atomics (sequential consistency)

Basic model for accelerators • One host device and multiple target devices of the same type. • Device = a logical execution engine • device data environment = a data environment associated with a target data or target region. • Keyword target constructs control how data and code is offloaded to a device. • Data is mapped from a host data environment to a device data environment.

Targets I • Code inside target region is executed on the device (default is sequential execution) • Parallel execution by other OpenMP directives • Clauses to control data movement. Mapping of variables • Can specify which device to use. TO the memory of device #pragma omp target map(to:B,C), map(tofrom:sum) For the result of reduction #pragma omp parallel for reduction(+:sum) for (int i=0;i
Targets II • target data construct just moves data and does not execute code – can have multiple target regions inside a target data region – allows data to persist on device between target regions

• target update construct updates data during a target data region. • declare target compiles a version of function/subroutine that can be called on the device.

Targets III • Target regions are blocking: the encountering thread waits for them to complete. – Asynchronous behaviour can be achieved by using target regions inside tasks (with dependencies if required).

• Executing a target region on a GPU can only use one CUDA SM because synchronization required for OpenMP – not possible between SMs = not too efficient

// OpenMP for loop parallelization void ompVectorAdd(int N, float *a,float *b,float *c) Mapping of variables { #pragma omp target map(to: N, a[0:N], b[0:N]) \ map(from: c[0:N]) { Output Inputs int i; #pragma omp parallel for for (i = 0; i < N; i++) c[i] = a[i] + b[i]; Parallel execution }}

Other functions Runtime support routines: • void omp_set_default_device(int dev_num) • int omp_get_default_device(void) • int omp_get_num_devices(void) • int omp_get_num_teams(void) • int omp_get_team_num(void) • int omp_is_initial_device(void) Environment variable: • Control default device through OMP_DEFAULT_DEVICE • Accepts a non-negative integer value

Map clause • Map a variable or an array section to a device data environment • Syntax: map(alloc | to | from | tofrom: list) • Map-types – alloc: allocate storage for corresponding variable – to: alloc and assign value of original variable to corresponding variable on entry – from: alloc and assign value of corresponding variable to original variable on exit – tofrom:default, both to and from

Terminology • Mapped variable: – original variable in a (host) data environment with a corresponding variable in a device data environment

• Mappable type: – A type that is amenable for mapped variables.(Bitwise copyable plus additional restrictions.)

Optimize data transfers • The target data construct creates a scoped device data environment – The map clauses control direction of data flow – The variables remain in the device data environment during the target data region

• Use target update to request data transfers from within a target data region

Optimize execution and data transfers #pragma omp target data device(0) map(alloc:tmp[:N]) map(to:input[:N)) map(from:res) { #pragma omp target device(0) #pragma omp parallel for for (i=0; i
Overlap of host and device code

extern void init(float*, float*, int); extern void init_again(float*, float*, int); extern void output(float*, int) void vec_mult(float *p, float *v1, float *v2, int N) { int i; init(v1, v2, N); #pragma omp target data map(from: p[0:N]) { #pragma omp target map(to: v1[:N], v2[:N]) #pragma omp parallel for for (i=0; i
Copy only once

Repeated copying

#pragma omp target data device(0) map(alloc:tmp[:N]) map(to:input[:N)) map(from:res) { #pragma omp target device(0) #pragma omp parallel for for (i=0; i
Update due to host modification

Advanced constructs • teams construct creates multiple master threads which can execute in parallel, and spawn parallel regions, but cannot – synchronize – or communicate with each other.

• distribute construct spreads the iterations of a parallel loop across teams. – only schedule option is static (with optional chunksize).

Asynchronous Offloading #pragma omp parallel sections { #pragma omp task { #pragma omp target map(to:input[:N]) map(from:result[:N]) #pragma omp parallel for for (i=0; i
Overlap of host and device code

Terminology • League – the set of threads teams created by a teams construct

• Contention group – threads of a team in a league and their descendant threads

“teams” constructs I • teams creates a league of thread teams – The master thread of each team executes the teams region – Number of teams is specified with num_teams() – Each team executes with thread_limit() threads

“teams” constructs I • teams constructs must be “perfectly” nested in a target construct: – No statements or directives outside the teams construct – Teams cannot synchronize

• Only special OpenMP constructs can be nested inside a teams construct: – – – –

distribute parallel parallel for parallel sections

SAXPY host version void saxpy(float *restrict y, float * restrict x, float a, int n) { #pragma omp parallel for for (int i = 0; i < n; ++i){ y[i] = a*x[i] + y[i]; }}

SAXPY device version I void saxpy(float *restrict y, float * restrict x, float a, int n) { #pragma omp target map(to:x[0:n], n, a) map(y[0:n]) { #pragma omp parallel for Simple version, due to synchronization it can use only one SM for (int i = 0; i < n; ++i){ y[i] = a*x[i] + y[i]; }}}

SAXPY device version II void saxpy(float * restrict y, float * restrict x, float a, int n) { int num_blocks = Fb(n); int nthreads = Ft(n); #pragma omp target data map(to:x[0:n], n, a) map(y[0:n]) #pragma omp teams num_teams(num_blocks) thread_limit(nthreads) #pragma omp parallel for for (int i = 0; i < n; i += num_blocks){ Generate league (CUDA:grid): teams (CUDA: blocks of threads) and threads for (int j = i; j < i + num_blocks; j++) { (CUDA: threads) y[j] = a*x[j] + y[j]; }} }}

Distribute construct Distribute the iterations of the associated loops across the master threads of a teams construct • No implicit barrier at the end of the construct • dist_schedule(kind[, chunk_size]) • Specified scheduling kind must be static – Chunks are distributed in round-robin fashion of chunks with size chunk_size

• If no chunk size specified, chunks are of (almost) equal size; each team receives at least one chunk

SAXPY device version III void saxpy(float * restrict y, float * restrict x, float a, int n) { int num_blocks = Fb(n); int nthreads = Ft(n); #pragma omp target data map(to:x[0:n], n, a) map(y[0:n]) #pragma omp teams num_teams(num_blocks) thread_limit(nthreads){ #pragma omp distribute for (int i = 0; i < n;i += num_blocks){ Another version: use distribute = no barrier #pragma omp parallel for for (int j = i; j < i + num_blocks; j++) { y[j] = a*x[j] + y[j]; } } }}

SAXPY device version IV void saxpy(float * restrict y, float * restrict x, float a, int n){ int num_blocks = Fb(n); int nthreads = Ft(n); #pragma omp target map(to:x[0:n], n, a) map(y[0:n]) #pragma omp teams distribute parallel for \ num_teams(num_blocks) thread_limit(bsize) for (int i = 0; i < n; ++i) y[i] = a*x[i] + y[i]; Shortened version of SAXPY } device version III

Task Dependencies void blocked_cholesky( int NB, float A[NB][NB] ) { int i, j, k; for (k=0; k
Multi-device Example int num_dev = omp_get_num_devices(); int chunksz = length / num_dev; assert((length % num_dev) == 0); #pragma omp parallel sections firstprivate(chunksz,num_dev) { for (int dev = 0; dev< NUM_DEVICES; dev++) { #pragma omp task firstprivate(dev) { int lb = dev * chunksz; Different devices can be used int ub = (dev+1) * chunksz; #pragma omp target device(dev) map(in:y[lb:chunksz]) map(out:x[lb:chunksz]) { #pragma omp parallel for for (int i = lb; i < ub; i++) { x[i] = a * y[i]; }}}}}

If Clause Example #define THRESHOLD1 1000000 #define THRESHOLD2 1000 extern void init(float*, float*, int); extern void output(float*, int); void vec_mult(float *p, float *v1, float *v2, int N) { int i; init(v1, v2, N); Threshold for the device #pragma omp target if(N>THRESHOLD1) map(to: v1[0:N], v2[:N]) execution map(from: p[0:N]) #pragma omp parallel for if(N>THRESHOLD2) for (i=0; i
Array Sections Example void foo (int *A, int N){ int *p; #pragma omp target data map( A[:N]) { // implicit map(tofrom: A) for the pointer // A = storage allocated for array section on device p = &A[0]; #pragma omp target map( p[0:N/2] ) { A[N-1] = 0; p[0] = 0; }}}

OpenACC • OpenACC (Open Accelerator): nový (představen 2012) standard pro práci s akcelerátory: – – – –

GPU Nvidia GPU AMD/ATI Xeon Phi atd.

• Přístup podobný OpenMP => vyšší úroveň • Optimalizace je přenechána kompilátoru => efektivita ? • Větší podpora jen v komerčních produktech: – Portland Group (PGI), Accelerator Compiler – CAPS , HMPP Workbench – Cray Corporation, Complilation Environment

• Nekomerční podpora: – projekt accULL – Pracuje se na podpoře v rámci GNU GCC

160

Násobení matic ( pomocí OpenMP) #pragma parallel for numthreads(size*size) schedule(static,1) \ shared(A, B, C, size) private(column, row, position) collapse(2) for(row = 0; row < size; row++) { for (column = 0; column < size; column++) { pom=0; for(int position = 0; position < size; position++) { pom+ = MA[row][position] * MB[position][column];} MC[row][column] = pom; }}

161

Násobení matic (GPU kód) void kernel (int row, int column) { int position ; float pom=0.0; for(int position = 0; position < size; position++) { pom += MA[row][position] * MB[position][column]; MC[row][column]=pom;} void main() { Nakopíruj data (matice A a B) z hlavní paměti do paměti GPU. Spusť size*size vláken (instancí kernelu) s příslušnými parametry row a column. Nakopíruj data (matice C) z paměti GPU do hlavní paměti. }

Zajímavé je, že se „ztratily“ dva for cykly. 162

OpenACC příklad #pragma acc kernels for(row = 0; row < size; row++) { for (column = 0; column < size; column++) { pom= 0; for(int position = 0; position < size; position++) { pom + = A[row][position]* B[position][column];} C[row][column]=pom; }} 163

Základy OpenACC • Obdobné OpenMP • Úprava programu pomocí direktiv #pragma acc directive-name [clause [[,] clause]…] new-line • Existují i OpenACC příkazy

Proměnné prostředí • ACC_DEVICE_TYPE: Defaultní typ akcelerátoru Př. ACC_DEVICE_TYPE=NVIDIA. • ACC_DEVICE_NUM: číslo defaultně použitého akcelerátoru

Základní konstrukce I Základní konstrukcí je #pragma acc parallel specifikuje části, které se budou nahrávat/provádět (offload) na akcelerátoru (je zakázáno vyskočit ven nebo zvnějšku skočit dovnitř). #pragma acc parallel { ... #pragma acc loop for(...){ ...loop body... } ... }

Základní konstrukce II #pragma acc parallel loop for(i=0;i
Upřesnění paralelismu Obdobně jako v OpenMP je možno nastavit: • num_gangs, • num_workers, • vector_length vector_length může být: 1, 64, 128 (default pro CUDA), 256, 512, 1024

Základní direktivy • kernels: obdoba parallel (z historických důvodů), výsledek kompilace se může lišit • copyin(): kopíruje souvislý kus paměti do akcelerátoru (před kernelem). • copyout(): kopíruje souvislý kus paměti z akcelerátoru (po kernelu). • copy(): kopíruje souvislý kus paměti do (před kernelem) a z akcelerátoru (po kernelu).

Alokace • create: vytvoří (alokuje) proměnnou na akcelarátoru Př. create(a[0:size][0:size]) • Pokud je znovupoužita stejná proměnná, present (pro zabránění opětovného kopírování nebo alokace) • Alokace mimo kernel: float *a = (float*) acc_malloc(sizeof(float)*size*size); acc_free(a); V kernelu je ten ukazatel označen jako deviceptr(a)

Paralelismus SW • gang (CUDA: blok vláken), • worker (CUDA: warp) • vector (CUDA: vlákno) • Př. #pragma acc loop gang(16), vector(32) HW (model architektury) • Několik processing elements (PEs), • každý PE je vícevláknový, • Každé vlákno PE může vykonávat vektorové instrukce.

Další možnosti loop + datové závislosti #pragma acc loop independent Vypíná automatickou kontrolu datových závislostí v cyklu (iterace se mohou provádět v libovolném pořadí) Opakem je: #pragma acc loop seq reduction provedení binární redukce (jako v OpenMP) cache označená data jsou označena pro uložení v rychlé paměti (např. CUDA sdílená paměť)