Succesorul chipurilor de prelucrare grafică din jocurile aracade, procesoarul grafic, acronim GPU (Graphics Processing Unit), este un circuit electronic, specializat, în crearea și manipularea imaginilor trimise către un afișaj electronic (e.g. monitor).

Utilitatea lui s-a extins ulterior către probleme “embarrassingly parallel”, iar în astăzi sunt folosite la antrenarea antrenarea rețelelor neurale și minarea de criptomonede. Vorbim aici despre întrebuințarea unui GPGPU (General Purpose GPU), un procesor grafic cu o flexibilitate ridicată de programare, capabil de a rezolva și probleme generale.

În execuție, o arhitectură de tip GPU folosește paradigma SIMD (Single Instruction Multiple Data, vezi taxonomia lui Flynn). Acesta presupune:

• schimb rapid de context între thread-uri,
• planificarea în grupuri de thread-uri,
• și orientare către prelucrari masive de date.

Deci, unitățile de tip GPU sunt potrivite pentru paralelismul de date, adică pentru un flux intensiv computațional, cu puține decizii de control.

De obicei, termenul de GPGPU apare atunci când unitatea GPU este folosită ca și coprocesor matematic. Astăzi, majoritatea unităților de tip GPU sunt și GPGPU. Ampla folosire a acestora se datorează:

• diferențelor de putere de procesare brută dintre CPU și GPU (instrucțiuni/secundă)
• standardizarea de API-uri care ușurează munca programatorilor
• răspândirea aplicațiilor ce pot beneficia de pe urma paralelismului de tip SIMD
• cererii pe piața unităților computaționale destinate:
  • atât consumatorilor (PC, Smartphone, TV, etc.),
  • cât și mediilor industriale (Automotive, HPC etc).

Principalii producatori de core-uri IP (intellectual property) tip GPU sunt:

• Intel http://en.wikipedia.org/wiki/List_of_Intel_graphics_processing_units
• Nvidia http://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units
• Amd http://en.wikipedia.org/wiki/List_of_AMD_graphics_processing_units
• Apple https://en.wikipedia.org/wiki/Apple_silicon
• Imagination http://en.wikipedia.org/wiki/List_of_PowerVR_products
• Qualcomm http://en.wikipedia.org/wiki/Adreno
• Vivante http://en.wikipedia.org/wiki/Vivante_Corporation

Dacă un IP de GPU este integrat pe aceeași pastilă de siliciu a unui SoC (System-on-a-Chip), spunem că este un GPU integrat. Exemple de SoC-uri cu IP de GPU integrat includ procesoarele x86 Intel/AMD, cât și majoritatea SoC-urilor pentru dispozitive mobile bazate pe arhitectura ARM (ex. Qualcomm Snapdragon). Un GPU integrat împarte ierarhia de memorie cu alte IP-uri (exp. controllere PCIe/USB/SATA/ETH).

De altfel, un GPU dedicat (discrete GPU) presupune valorificarea unei unui spațiu de memorie, mapat peste VRAM (Video Random-Access Memory), cât și o magistrală PCIe/AGP8x/USB pentru comunicare cu sistemul. Exemple de GPU-uri dedicate sunt seriile de plăci grafice Geforce (Nvidia) și Radeon (AMD).

Exemple de folosire de GPGPU-uri: prelucrări video, audio și de imagini, simulări ale fenomenelor fizice, finanțe, criptografie, design electronic (VLSI), mașini autonome.

• https://www.nvidia.com/en-us/self-driving-cars/partners/bmw/
• https://blogs.nvidia.com/blog/2018/09/18/audi-unveils-e-tron-electric-suv/

Rețele neurale - antrenare vs. inferență.

• https://www.forbes.com/sites/forbestechcouncil/2017/12/01/for-machine-learning-its-all-about-gpus/
• https://www.quora.com/Why-are-GPUs-well-suited-to-deep-learning

Criptomonede - mining via hashing.

• https://coincentral.com/best-gpu-for-mining-2018-edition/

SmartTV, Smartphone - accelerare video, recunoaștere facială/audio.

Simulări fizice - NVIDIA Physx, Folding@Homel

• https://blogs.nvidia.com/blog/2018/11/13/weather-predicted-sc18-gpu-hpc-jensen-huang/

Prelucrări multimedia - filtre imagini GIMP/Photoshop.

Alte domenii - arhivare (WinZip), encriptare.

În cadrul unui sistem ce conține un GPU, procesorul general (CPU) coordonează execuția și este numit “HOST”; pe când unitatea care efectuează calculele (GPU) este numită “DEVICE”.

HOST-ul controlează toate schimbarile de stare în cadrul unui GPU, alocările/transferurile de memorie și evenimentele ce țin de sistemul de operare.

O unitate GPU conține un procesor de comandă (“command processor”) care citește comenzile scrise de către HOST (CPU) în anumite zone de memorie mapate spre access atât către unitatea GPU, cât și către CPU.

În cazul GPU-urilor dedicate, o prelucrare de date necesită în prealabil un transfer din RAM către VRAM. Acest transfer se face printr-o magistrală (PCIe, AGP, USB…). Viteza de transfer RAM-VRAM via magistrală este inferioară vitezei de acces la RAM sau la VRAM. O potențială optimizare în cadrul acestui transfer ar fi intercalarea cu procesarea.

În cazul GPU-urilor integrate transferul RAM↔“VRAM” presupune o mapare de memorie, de multe ori translatată printr-o operație de tip zero-copy.

Programarea unui GPU se face printr-un API (Application Programming Interface). Cele mai cunoscute API-uri orientate către folosirea unui GPU ca coprocesor matematic sunt: CUDA, OpenCL, DirectCompute, OpenACC și Vulkan.

Dezvoltarea de cod pentru laboratoarele de GPU se va face folosind CUDA.

CUDA este un API introdus în 2006 de Nvidia pentru GPU-urile sale. În prezent CUDA este standardul de facto pentru folosirea unităților GPU în industrie și cercetare. Aceasta se datorează faptului că este o platformă stabilă cu multe facilități. Dacă o nouă versiune de CUDA introduce noi funcționalități, dar arhitectura nu le suportă, acestea sunt dezactivate.

În mare toate GPU-urile oferite de Nvidia sunt suportate, diferența fiind la facilitățile suportate. Singura limitare, majoră, a platformei CUDA este că suportă numai unități de procesare de tip GPU de la Nvidia.

Un standard alternativ la CUDA este OpenCL, suportat de Khronos și implementat de majoritatea producătorilor de GPU (inclusiv Nvidia ca o extensie la CUDA). OpenCL suferă de următoarele lipsuri: * suportul este fragmentat * standardul este mult mai restrictiv (decât CUDA) * mai complicat de scris programe (decât CUDA)

CUDA (Compute Unified Device Architecture) permite utilizarea limbajului C pentru programarea pe GPU-urile Nvidia cât și extensii pentru alte limbaje (exp. Python). Deoarece una din zonele țintă pentru CUDA este HPC (High Performance Computing), în care limbajul Fortran este foarte popular, PGI ofera un compilator de Fortran ce permite generarea de cod și pentru GPU-urile Nvidia. Există binding-uri pentru Java (jCuda), Python (PyCUDA) și .NET (CUDA.NET).

Unitatea de bază în cadrul arhitecturii CUDA este numită SM (Streaming Multiprocessor). Ea conține în funcție de generație un număr variabil de CUDA Cores sau SP (Stream Processors) - de regulă între 8SP și 128SP. Unitatea de bază în scheduling este denumită “warp” și este alcatuită din 32 de thread-uri. Vom aborda mai amănunțit arhitectura CUDA în laboratorul următor.

Versiunea de "compute capability" a unui SM, are formatul X.Y, unde X este versiunea majoră, pe când Y este versiunea minoră. Partea majoră identifică generația din care face parte arhitectura.

Partea minoră identifică diferențe incrementale în arhitectură și posibile noi funcționalități.

Știind versiunea majoră și cea minoră cunoaștem facilitățile hardware oferite de către arhitectură.

O listă a GPU-urile NVIDIA și versiunile lor majore/minore se regăsește aici.

CUDA extinde limbajul C prin faptul că permite unui programator să definească funcții C, denumite kernels, care urmează a fi execute în paralel de N thread-uri CUDA. Scopul este de a abstractiza arhitectura GPU astfel încat partea de scheduling cât și gestiunea resurselor se face de catre stiva software CUDA împreună cu suportul hardware. Figura de mai jos denotă distribuirea thread-urilor către două arhitecturi partiționate diferit.

Un kernel se definește folosind specificatorul global iar execuția sa se face printr-o configurație de execuție folosind <<<...>>> . Configurația de execuție denotă numarul de blocuri și numărul de thread-uri dintr-un block. Fiecare thread astfel poate fi identificat unic prin blockIdx și threadIdx.

Mai jos avem definit un kernel, vector_add, care are ca argumente pointeri de tip float, respectiv size_t. Acesta calculează $ f(x) = 2x + 1/(x + 1) $, pentru fiecare element din vector. Numărul total de thread-uri este dimensiunea vectorului.

__global__ void vector_add(const float *a, float *b, const size_t n) {
   // Compute the global element index this thread should process
   unsigned int i = threadIdx.x + blockDim.x * blockIdx.x;
 
   // Avoid accessing out of bounds elements
   if (i < n) {
      b[i] = 2.0 * a[i] + 1.0 / (a[i] + 1.0);
   }
}

Configurația de execuție denotă maparea între date și instrucțiuni. În funcția de kernel, se definește setul de instrucțiuni ce se va executa repetat pe date. Mai jos vector_add este lansat în execuție cu N thread-uri (blocks_no x block_size) organizate câte block_size thread-uri per bloc.

// Launch the kernel
vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);

#include <stdio.h>
 
__global__ void kernel_example(int value) {
   /**
   * This is a kernel; a kernel is a piece of code that
   * will be executed by each thread from each block in
   * the GPU device.
   */
   printf("[GPU] Hello from the GPU!\n");
   printf("[GPU] The value is %d\n", value);
   printf("[GPU] blockDim=%d, blockId=%d, threadIdx=%d\n",blockDim.x, blockIdx.x, threadIdx.x);
}
 
int main(void) {
   /**
   * Here, we declare and/or initialize different values or we
   * can call different functions (as in every C/C++ program);
   * In our case, here we also initialize the buffers, copy
   * local data to the device buffers, etc (you'll see more about
   * this in the following exercises).
   */
   int nDevices;
   printf("[HOST] Hello from the host!\n");
 
   /**
   * Get the number of compute-capable devices. See more info 
   * about this function in the Cuda Toolkit Documentation.
   */
   cudaGetDeviceCount(&nDevices);
   printf("[HOST] You have %d CUDA-capable GPU(s)\n", nDevices);
 
   /** 
   * Launching the above kernel with a single block, each block
   * with a single thread. The syncrhonize and the checking functions
   * assures that everything works as expected.
   */
   kernel_example<<<1,1>>>(25);
   cudaDeviceSynchronize();
 
   /**
   * Here we can also deallocate the allocated memory for the device
   */
   return 0;
}

O aplicatie CUDA are ca scop executia de cod pe GPU-uri NVIDIA CUDA. In cadrul laboratoarelor partea de CPU (host) va fi folosita exclusiv pentru managementul executiei partii de GPU (device). Aplicatiilor vor viza executia folosind un singur GPU NVIDIA CUDA.

In codul prezentat mai jos, functia vector_add este marcata cu “global” si va fi compilata de catre CUDA NVCC compiler pentru GPU-ul de pe sistem (in cazul cozii xl va fi NVIDIA Pascal P100).

/**
 * This kernel computes the function f(x) = 2x + 1/(x + 1) for each
 * element in the given array.
 */
__global__ void vector_add(const float *a, float *b, const size_t n) {
  // Compute the global element index this thread should process
  unsigned int i = threadIdx.x + blockDim.x * blockIdx.x;
 
  // Avoid accessing out of bounds elements
  if (i < n) {
    b[i] = 2.0 * a[i] + 1.0 / (a[i] + 1.0);
  }
}

Din punct de vedere hardware, partea de host (CPU) are ca memorie principala RAM (chip-uri memorie instalate pe placa de baza via slot-uri memorie) iar partea de device (GPU) are VRAM (chip-uri de memorie prezente pe placa video). Cand vorbim de memoria host (CPU) ne referim la RAM, iar in cazul memoriei device (GPU) la VRAM.

La versiunile mai recente de CUDA, folosind limbajul C/C++, un pointer face referire la spatiul virtual care este unificat pentru host (CPU) si device (GPU). Adresele virtuale insa sunt translatate catre adrese fizice ce rezida ori in memoria RAM (CPU) ori in memoria VRAM (GPU). Astfel este important cum alocam memoria (fie cu malloc pentru CPU sau cudaMalloc pentru GPU) si respectiv sa facem cu atentie transferurile de memorie intre zonele virtuale definite (de la CPU la GPU si respectiv de la GPU la CPU).

// Declare variable to represent ~1M float values and
// computes the amount of bytes necessary to store them
const int num_elements = 1 << 16;
const int num_bytes = num_elements * sizeof(float);
 
// Declaring the 'host arrays': a host array is the classical
// array (static or dynamically allocated) we worked before.
float *host_array_a = 0;
float *host_array_b = 0;
 
// Declaring the 'device array': this array is the equivalent
// of classical array from C, but specially designed for the GPU
// devices; we declare it in the same manner, but the allocation
// process is going to be different
float *device_array_a = 0;
float *device_array_b = 0;

Functia malloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din RAM.

// Allocating the host array
host_array_a = (float *) malloc(num_bytes);
host_array_b = (float *) malloc(num_bytes);

Functia cudaMalloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din VRAM.

// Allocating the device's array; notice that we use a special
// function named cudaMalloc that takes the reference of the
// pointer declared above and the number of bytes.
cudaMalloc((void **) &device_array_a, num_bytes);
cudaMalloc((void **) &device_array_b, num_bytes);
 
// If any memory allocation failed, report an error message
if (host_array_a == 0 || host_array_b == 0|| device_array_a == 0 || device_array_b == 0) {
  printf("[HOST] Couldn't allocate memory\n");
  return 1;
}

// Initialize the host array by populating it with float values  
for (int i = 0; i < num_elements; ++i) {
  host_array_a[i] = (float) i;
}
 
// Copying the host array to the device memory space; notice the
// parameters of the cudaMemcpy function; the function default
// signature is cudaMemcpy(dest, src, bytes, flag) where
// the flag specifies the transfer type.
//
// host -> device: cudaMemcpyHostToDevice
// device -> host: cudaMemcpyDeviceToHost
// device -> device: cudaMemcpyDeviceToDevice
cudaMemcpy(device_array_a, host_array_a, num_bytes, cudaMemcpyHostToDevice);

// Compute the parameters necessary to run the kernel: the number
// of blocks and the number of threads per block; also, deal with
// a possible partial final block
const size_t block_size = 256;
size_t blocks_no = num_elements / block_size;
 
if (num_elements % block_size) 
  ++blocks_no;
 
// Launch the kernel
vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
cudaDeviceSynchronize();

// Copy the result back to the host memory space
cudaMemcpy(host_array_b, device_array_b, num_bytes, cudaMemcpyDeviceToHost);
 
// Print out the first 10 results
for (int i = 0; i < 10; ++i) {
  printf("Result %d: 2 * %1.1f + 1.0/(%1.1f + 1.0)= %1.3f\n", 
  i, host_array_a[i], host_array_a[i], host_array_b[i]);
}

// Deallocate memory
free(host_array_a);
free(host_array_b);
cudaFree(device_array_a);
cudaFree(device_array_b);

mkdir asc
cd asc
pwd # acesta va returna <cale_absolută_director_asc_fep8>

mkdir asc # puteți da comanda din directorul vostru home
cd asc
pwd # acesta va returna <cale_absolută_director_asc_local>
sudo chown -R <user> asc # numele user-ului vostru de pe local
decomentăm linia 'user_allow_other' din /etc/fuse.conf
sshfs -o allow_other <user.moodle>@fep8.grid.pub.ro:<cale_absolută_director_asc_fep8> <cale_absolută_director_asc_local>

Urmăriți instrucțiunile de pe GitLab GitLab.

Enunt Laborator 4

• Responsabili laborator: Costin Carabaș, Tudor Calafeteanu, Grigore Lupescu, Mihnea Mitroi, Irinel Gul, Alex Bala

• Documentatie CUDA:
  • CUDA C Programming
  • CUDA NVCC compiler
  • CUDA Visual Profiler
  • CUDA Dev Toolkit
  • CUDA GPUs
• Acceleratoare xl (NVidia P100)
  • NVIDIA Pascal P100
• Advanced CUDA
  • CUDA Streams
  • CUDA Thread Basics
  • An Even Easier Introduction to CUDA