Programarea de kernel in OpenCL

In laboratorul anterior am introdus notiunea de kernel anume codul ce urmeaza sa fie executat de catre DEVICE.

Cum OpenCL defineste atat un API pentru partea de HOST cat si pentru partea DEVICE (kernel) acestea pot sa fie diferite ca versiune. Mai mult, pe un sistem pot fi instalate mai multe implementari de OpenCL si este important de avut in vedere ca 2 implementari de HOST pot avea acces la acelasi DEVICE.

Un exemplu ipotetic ar fi un sistem cu 3 DEVICES: 1xCPU Intel Xeon si 2xGPUs NVIDIA Tesla. Astfel am avea
→ PLATFORM_0: AMD OpenCL, OpenCL 2.1 → DEVICE_0: Intel Xeon CPU, OpenCL 1.2
→ PLATFORM_1: INTEL OpenCL, OpenCL 2.0 → DEVICE_0: Intel Xeon CPU, OpenCL 2.0
→ PLATFORM_2: NVIDIA CUDA, OpenCL 1.1 → DEVICE_0: NVIDIA Tesla, OpenCL 1.0, DEVICE_1: NVIDIA Tesla, OpenCL 1.0

In acest caz DEVICE_0 este vazut de partea HOST atat in PLATFORM_0 cat si in PLATFORM_1 insa reprezinta acelasi device, anume Intel Xeon CPU. In exemplul de mai sus, folosind libOpenCL de la AMD vom avea suport pe partea HOST de OpenCL 2.1 si pe partea de DEVICE OpenCL 1.2 pe cand folosind varianta de la INTEL va oferi suport OpenCL 2.0 (HOST) / OpenCL 2.0 (DEVICE).

In particular pe sistemele ibm-dp.q sau hp-sl.q avem disponibile 2 platforme, NVIDIA si INTEL cu urmatoarele biblioteci:
NVIDIA → /export/opt/tools/intel/opencl-1.2-6.4.0.25/lib64/libintelocl.so
INTEL → /usr/local/cuda-9.0/targets/x86_64-linux/lib/libOpenCL.so

Pentru programarea de kernel ne intereseaza versiunea de OpenCL pentru DEVICE, anume OpenCL 1.1 pentru NVIDIA Tesla 2050 sau OpenCL 1.2 pentru NVIDIA Tesla K40M. Mai multe detalii despre versiunile de OpenCL de pe cluster se regasesc aici ibm-dp.q, ibm-dp48.q si aici hp-sl.q.

Ultima versiune OpenCL definita de Khronos este OpenCL 2.2 insa versiunile implementate de vendori sunt OpenCL 1.1, OpenCL 1.2 sau OpenCL 2.0. Pe cluster am putea sa ne limitam la OpenCL 1.1 pentru a putea rula atat pe ibm-dp.q cat si pe hp-sl.q.

Imaginea de mai jos prezinta o mapare a problemei (dimensiune globala 3D), alcatuit din mai multe workgroups. La randul sau un workgroup (dimensiune locala 3D) este alcatuit din mai multe work-items.

Este foarte important de avut in vedere caracteristicile arhitecturii unde va rula codul de kernel. Astfel maparea din cauza lipsei de resurse (ex registri, latime banda memorie) poate fi foarte ineficienta ducand la o utilizare foarte scazuta a hardware-ului.

Devices CLUSTER

Pe cluster avem urmatoarele cozi: ibm-dp.q, ibm-dp48.q si hp-sl.q.

Device GPU Tesla K40m (hp-sl.q), API OpenCL, platforma NVIDIA:

  Device Name                                     Tesla K40m
  Device OpenCL C Version                         OpenCL C 1.2
  Device Type                                     GPU
  Device Profile                                  FULL_PROFILE
  Device Topology (NV)                            PCI-E, 14:00.0
  Max compute units                               15
  Max clock frequency                             745MHz
  Compute Capability (NV)                         3.5
  Max work item dimensions                        3
  Max work item sizes                             1024x1024x64
  Max work group size                             1024
  Preferred work group size multiple              32
  Warp size (NV)                                  32

Device GPU Tesla M2070 (ibm-dp.q), API OpenCL, platforma NVIDIA:

  Device Name                                     Tesla M2070
  Device OpenCL C Version                         OpenCL C 1.1 
  Device Type                                     GPU
  Device Profile                                  FULL_PROFILE
  Device Topology (NV)                            PCI-E, 14:00.0
  Max compute units                               14
  Max clock frequency                             1147MHz
  Compute Capability (NV)                         2.0
  Max work item dimensions                        3
  Max work item sizes                             1024x1024x64
  Max work group size                             1024
  Preferred work group size multiple              32
  Warp size (NV)                                  32

Device CPU Xeon E5-2670 (hp-sl.q), API OpenCL, platforma INTEL:

  Device Name                                     Intel(R) Xeon(R) CPU E5-2670 v2@ 2.50GHz
  Device OpenCL C Version                         OpenCL C 1.2
  Device Type                                     CPU
  Device Profile                                  FULL_PROFILE
  Max compute units                               40
  Max clock frequency                             2500mhz
  Max work item dimensions                        3
  Max work item sizes                             8192x8192x8192
  Max work group size                             8192
  Compiler Available                              Yes
  Linker Available                                Yes
  Preferred work group size multiple              128

Device CPU Xeon X5650 (ibm-dp.q), API OpenCL, platforma INTEL:

  Device Name                                     Intel(R) Xeon(R) CPU X5650  @ 2.67GHz
  Device OpenCL C Version                         OpenCL C 1.2
  Device Type                                     CPU
  Device Profile                                  FULL_PROFILE
  Max compute units                               24
  Max clock frequency                             2670MHz
  Max work item dimensions                        3
  Max work item sizes                             8192x8192x8192
  Max work group size                             8192
  Compiler Available                              Yes
  Linker Available                                Yes
  Preferred work group size multiple              128

Arhitectura GPU NVIDIA CUDA

Laboratorul 9 discuta arhitectura de GPU mai in detaliu. Ideal de parcurs rapid si acesta ca o completare.

Arhitectura NVIDIA FERMI aici, Tesla 2050, fep queue ibm-dp.q

Arhitectura NVIDIA KEPLER aici, Tesla K40M, fep queue hp-sl.q

Prima arhitectura NVIDIA complet programabila a fost G80 (ex. Geforce 8800, lansat in anul 2006). Cu aceasta arhitectura s-a trecut de la unitati hardware fixe vertex/pixel la cele de unified shader care puteau procesa atat vertex/pixel cat si geometry. Evolutia arhitecturilor GPU de la NVIDIA este detaliata aici.

Implementarea NVIDIA pentru GPGPU se numeste CUDA (Compute Unified Device Architecture) si permite utilizarea limbajului C pentru programarea pe GPU-urile proprii. Lista de GPU-uri ce suporta API-ul CUDA sau OpenCL se regaseste pe site-ul oficial aici sau pe wiki aici. Fiecare noua arhitectura are un codename (ex Fermi, Pascal) si este reprezentata de un “compute capability” (list aici). Cu cat arhitectura este mai noua, cu atat sunt suportate mai multe facilitati din API-urile CUDA si OpenCL.

Framework-ul CUDA implementeaza ca extensie si API-ul de OpenCL prin intermediul caruia vom interactiona cu GPGPU-ul NVIDIA Tesla disponibil pe ibm-dp.q sau hp-sl.q. Din pacate API-ul de OpenCL are suport slab fata de API-ul propriu CUDA. Multe din programele care fac debug/profiling nu au suport in OpenCL sau au avut suportul deprecated sau eliminat precum NVIDIA Visual Profiler care era suportat intr-o versiune anterioara de CUDA.

Unitatea GPU este potrivita pentru paralelismul de date SIMD (Single Instruction Multiple Data), astfel aceleasi instructiuni sunt executate in paralel pe mai multe unitati de procesare. Datorita faptului ca acelasi program este executat pentru fiecare element de date, sunt necesare mai putine elemente pentru controlul fluxului. Si deoarece calculele sunt intensive computational, latenta accesului la memorie poate fi ascunsa prin calcule in locul unor cache-uri mari pentru date.

Motivul discrepantei intre performanta paralela dintre CPU si GPU este faptul ca GPU sunt specializate pentru procesare masiv paralela si intensiva computational (descrierea perfecta a taskurilor de randare grafica) si construite in asa fel incat majoritatea tranzistorilor de pe chip se ocupa de procesarea datelor in loc de cachingul datelor si controlul fluxului executiei.

La GPU-urile NVIDIA, un Streaming Processor (SP) este un microprocesor cu executie secventiala, ce contine un pipeline, unitati aritmetico-logice (ALU) si de calcul in virgula mobila (FPU). Nu are un cache, fiind bun doar la executia multor operatii matematice. Un singur SP nu are performante remarcabile, insa prin cresterea numarului de unitati, se pot rula algoritmi ce se preteaza paralelizarii masive.

SP impreuna cu Special Function Units (SFU) sunt incapsulate intr-un Streaming Multiprocessor (SM/SMX). Fiecare SFU contine unitati pentru inmultire in virgula mobila, utilizate pentru operatii transcendente (sin, cos) si interpolare. MT se ocupa cu trimiterea instructiunilor pentru executie la SP si SFU.

Pe langa acestea, exista si un cache (de dimensiuni reduse) pentru instructiuni, unul pentru date precum si memorie shared, partajata de SP-uri. Urmatorul nivel de incapsulare este Texture / Processor Cluster (TPC). Acesta contine SM-uri, logica de control si un bloc de handling pentru texturi. Acest bloc se ocupa de modul de adresare al texturilor, logica de filtrare a acestora precum si un cache pentru texturi.

Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru este facut posibil prin paralelismul existent la nivel hardware.

Documentatia NVIDIA recomanda rularea unui numar cat mai mare threaduri pentru a executa un task. Arhitectura CUDA de exemplu suporta zeci de mii de threaduri, numarul acestora fiind mult mai mare decat unitatile fizice existente pe chip. Acest lucru se datoreaza faptului ca un numar mare de threaduri poate masca latenta accesului la memorie.

Urmarind acelasi model modular ca si arhitectura, threadurile sunt incapsulate in blocuri (thread blocks / warps), iar blocurile in grile (thread grid). Fiecare thread este identificat prin indexul threadului in bloc, indexul blocului in grila si indexul grilei. Indexurile threadurilor si ale blocurilor pot fi uni/bi/tri-dimensionale, iar indexul grilei poate fi uni sau bi-dimensional. Acest tip de impartire are rolul de a usura programare pentru probleme ce utilizeaza structuri de date cu mai multe dimensiuni. Se poate observa ca thread-urile dintr-un thread block trebuie sa execute cat mai multe instructiuni identice spre a nu irosi resurse.

{{thread.blocks.jpg|''Structura threadurilor in blocuri''

Threadurile dintr-un bloc pot coopera prin partajarea de date prin intermediul memoriei shared si prin sincronizarea executiei. Functia de bariera functioneaza doar pentru threadurile dintr-un bloc. Sincronizarea nu este posibila la alt nivel (intre blocuri/grila etc.). Mai multe explicatii se regasesc in Laboratorul 9.

OpenCL date scalare si vectoriale

Pentru a transpune si utiliza eficient registri, OpenCL dispune atat de date tip scalar cat si de date tip vectorial.
Arhitectura interna denota utilitatea acestor tipuri de date. Pentru arhitecturile noi de tip GPU nu exista avantaje in a folosi tipuri de date vectoriale. Arhitecturile mai vechi de GPU bazate pe VLIW (Very Long Instruction Word) cat si arhitecturi tip CPU (SSE, AVX, NEON) beneficiaza de folosirea unor date de tip vectorial. Exemplu de folosire a datelor tip vectorial in OpenCL aici

Exemple date tip scalar:

OpenCL Type (device) API Type (host) Descriere
char cl_char 8 bit
int cl_int 32 bit
float cl_float 32 bit

Exemple date tip vectorial:

Tipul de date Descriere Valori
charn cl_charn 8 bit
intn cl_intn 32 bit
floatn cl_floatn 32 bit

In codul urmator definim 3 variabile de tip vectorial int4 si facem suma.

int4 vA = (int4)(7, -3, -2, 5);
int4 vB = (int4)(1, 2, 3, 4);
int4 vSum = vA + vB; // vSum va fi (8, -1, 1, 9)

In codul urmator definim 2 variabile de tip vectorial float4 si facem inmultirea cu un scalar.

float4 vf = (float4)(3.0f, -1.0f, 1.0f, -2.0f);
float4 vMul = vf * 2; // vMul va fi (6.0f, -2.0f, 2.0f, -4.0f)

In cazul arhitecturii NVIDIA TESLA (GPU) este de preferat folosirea tipurilor de date scalare. Paralelismul este exclusiv la nivel de work-item. Informatii OpenCL Tesla 2070:

  Preferred / native vector sizes
    char               1 / 1
    short              1 / 1
    int                1 / 1
    long               1 / 1
    half               0 / 0        (n/a)
    float              1 / 1
    double             1 / 1        (cl_khr_fp64)

Utilizarea variabilelor de tip vectorial poate reduce performanta in cazul NVIDIA Tesla.

In cazul arhitecturii INTEL XEON (CPU) este de preferat folosirea tipurilor de date vectoriale. Paralelismul este astfel atat la nivel de work-item (thread) cat si la nivel de operatii cu date vectoriale (via SSE, AVX). Informatii OpenCL Intel Xeon 5650:

  Preferred / native vector sizes
    char               1 / 16
    short              1 / 8
    int                1 / 4
    long               1 / 2
    half               0 / 0        (n/a)
    float              1 / 4
    double             1 / 2        (cl_khr_fp64)

Neutilizarea variabilelor de tip vectorial poate reduce performanta in cazul INTEL Xeon. Este posibil astfel ca unitatile SSE/AVX din CPU, sa nu fie folosite la capacitatea maxima.

OpenCL operatii in virgula mobila

Standardul OpenCL suporta urmatoarele precizii de calcul in virgula mobila:
HALF → 16 bit
FLOAT → 32 bit
DOUBLE → 64 bit

Calculul in dubla precizie (double IEEE 754) in OpenCL este optional, dar suportat in cazul NVIDIA Tesla M2070/K40M cat si INTEL XEON 5650/E5-2670. Pentru activare trebuie definit explicit in fisierul de kernel OpenCL. Operatiile de tip double (64 bit) ofera o precizie mult mai buna fata de cele de tip float (32 bit) insa performanta e considerabil mai redusa. Spre exemplu aici avem un tabel cu performanta mai multor arhitecturi Tesla atat in precizie single si double.

In unele cazuri insa, precizia double ar putea fi prea mult si ar fi de preferat precizia single sau chiar precizie half daca aceasta ofera o performanta mai buna. Un exemplu de astfel de problema sunt retelele neurale.

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

Caracteristici relevante operatii in virgula mobila API OpenCL, platforma NVIDIA, device Tesla M2070:

  Half-precision Floating-point support           (n/a)
  Single-precision Floating-point support         (core)
    Denormals                                     Yes
    Infinity and NANs                             Yes
    Round to nearest                              Yes
    Round to zero                                 Yes
    Round to infinity                             Yes
    IEEE754-2008 fused multiply-add               Yes
    Correctly-rounded divide and sqrt operations  Yes
  Double-precision Floating-point support         (cl_khr_fp64)
    Denormals                                     Yes
    Infinity and NANs                             Yes
    Round to nearest                              Yes
    Round to zero                                 Yes
    Round to infinity                             Yes
    IEEE754-2008 fused multiply-add               Yes

Functiile matematice primesc ca argumente ori date de tip scalar ori de tip vectorial.
Exemple functii matematice generale :

Functie
T acos(T) Arc cosine
T asin(T) Arc sine
T exp(T) Exponential base e
T exp2(T) Exponential base 2
T exp10(T) Exponential base 10
T mad(Ta, Tb, Tc) Approximates a * b + c

Functiile matematice pot avea suport integral hardware sau doar in sofware. Prin compilarea de kernel cu anumite flaguri precum “-cl-fast-relaxed-math” putem imbunatati performanta cu pretul preciziei. La baza, unitatile GPU au avut implementate functii matematice cu precizie variabila spre a avea o performanta cat mai mare. Seria de GPGPU-uri Tesla este orientata strict spre calcul tip HPC si astfel are ajustata arhitectura spre a integra unitati de inalta precizie atat pentru float cat si double (IEEE 754).

Profiling in OpenCL

Pentru a evalua timpul unei operatii trimise in coada de executie avem nevoie de suport special atat hardware cat si software. In acest moment NVIDIA nu ofera unelte software dedicate pentru OpenCL ci doar pentru CUDA. In API-ul de OpenCL sunt definite insa metode a putea face profiling pe care stack-ul celor de la NVIDIA le implementeaza.

Daca se doreste folosirea proprietatilor de profiling, trebuie creata coada de executie (command queue) cu proprietatea CL_QUEUE_PROFILING_ENABLE.

cmdQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &ret);

Interogarea proprietatilor OpenCL referitor la coada de executie, ibm-dp.q, platforma NVIDIA, device Tesla M2070:

  Queue properties
    Out-of-order execution                        Yes
    Profiling                                     Yes
  Profiling timer resolution                      1000ns
  Execution capabilities
    Run OpenCL kernels                            Yes
    Run native kernels                            No
    Kernel execution timeout (NV)                 No
  Concurrent copy and kernel execution (NV)       Yes
    Number of async copy engines                  2

Interogarea proprietatilor OpenCL referitor la coada de executie, ibm-dp.q, platforma INTEL, device Xeon X5650:

  Queue properties
    Out-of-order execution                        Yes
    Profiling                                     Yes
    Local thread execution (Intel)                Yes
  Prefer user sync for interop                    No
  Profiling timer resolution                      1ns
  Execution capabilities
    Run OpenCL kernels                            Yes
    Run native kernels                            Yes
    SPIR versions                                 1.2
  printf() buffer size                            1048576 (1024KiB)

Functiile de profiling definite in OpenCL API ofera informatii cu privire la executia de kernel/comenzi pe DEVICE spre deosebire de functii precum clock care nu poate masura decat timpul pe HOST.

Timpul de executie al unui kernel vazut pe partea de HOST (clEnqueueKernel+clFinish) este diferit de cel efectiv pe partea de DEVICE.

Pentru a realiza profiling folosind API-ul de OpenCL se folosesc evenimente la executia unei operatii pe coada (cmdQueue). Se defineste o variabila eveniment de tip cl_event care apoi se foloseste in functii precum clGetEventProfilingInfo, respectiv clWaitForEvents.
Mai jos este prezentat un exemplu de profiling folosind API-ul de OpenCL aici.

g_cmd_queue = clCreateCommandQueue(…CL_QUEUE_PROFILING_ENABLE, NULL);
clEnqueueNDRangeKernel(g_cmd_queue,…, &perf_event);
clWaitForEvents(1, &perf_event);
cl_ulong start = 0, end = 0;
clGetEventProfilingInfo(perf_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
clGetEventProfilingInfo(perf_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
 
//END-START gives you hints on kind of “pure HW execution time”
//the resolution of the events is 1e-09 sec
g_NDRangePureExecTimeMs = (cl_double)(end - start)*(cl_double)(1e-06);

Aplicatii dedicate profiling OpenCL

Exista mai multe aplicatii de profiling runtime pentru OpenCL dintre care AMD Code XL, Nvidia Visual Profiler (<CUDA 6.5, deprecated acum) sau Intel Platform Analyzer. Versiunile de OpenCL instalate pe fep permit doar profiling prin API-ul de OpenCL.

Exercitii

Urmariti indicatiile todo si documentatia oficiala OpenCL pentru a rezolva exercitiile.
OpenCL refcard 1.1 prezinta o listare a tuturor functiilor OpenCL 1.1 - atat host cat si device.

Intrati pe frontend-ul fep.grid.pub.ro folosind contul de pe cs.curs.pub.ro. Executați comanda qlogin -q ibm-dp.q pentru a accesa una din stațiile cu GPU-uri. Modificarile se vor face in host.cpp (ex 1,2) cat si in device.cl (ex 2) - urmariti indicatiile TODO din cod.

  1. (2p) Completati kernelul astfel incat sa se efectueze inmultirea de matrice C=AxB
  2. (2p) Completati fisierul “host.cpp” astfel incat sa faceti profiling la executia de kernel. Adaugati atributul de profiling la coada de executie (command queue). Afisati timpul de executie al functiei kernel folosind profiling cu events.
  3. (2p) Efectuati profiling la transferurile de memorie (ex. clEnqueueReadBuffer).
  4. (1p) Adaugati la kernel optiuni de compilare (ex. “-cl-fast-relaxed-math”)
  5. (2p) Modificati functia de kernel astfel incat sa folositi tipuri de date vectoriale (ex. float4). Exemplu aici. Verificati performanta arhitecturii GPU Tesla in comparatie cu arhitectura CPU Xeon relativ la folosirea datelor tip scalar si vectorial. Ce observati?
  6. (1p) Descarcati si compilati sursele aplicatiei clinfo (link se afla mai jos in sectiunea de referinte). Listati intreg setul de proprietati al device-ul NVIDIA Tesla (GPU) si respectiv al procesoarelor Intel Xeon (CPU). Selectati cele mai importante caracteristici ale celor doua arhitecturi si comentati-le.

Resurse

Schelet Laborator 8

Solutie Laborator 8

Enunt Laborator 8

  • Responsabil laborator: Grigore Lupescu

Referinte

asc/lab8/index.txt · Last modified: 2018/09/02 13:43 by grigore.lupescu
CC Attribution-Share Alike 3.0 Unported
www.chimeric.de Valid CSS Driven by DokuWiki do yourself a favour and use a real browser - get firefox!! Recent changes RSS feed Valid XHTML 1.0