Intel Xeon Phi

Az oldalon található információk segítséget nyújtanak a HPC felhasználóinknak, hogy alkalmazásaikat miként portolhatják valamint optimalizálhatják a hatékonyabb futást szem előtt tartva Intel Xeon Phi koprocesszor kártyákon.

Hardver bemutatása

Intel® Xeon Phi™ Coprocessor 7120P

• 

CPU	threads/Core	max threads	CPU Freq	Memória	L1 cache	L2 cache	Interfész	Ptot	Számítási kapacitás
61	4 Hardver	244	1238 GHz	16GB GDDR5	32KB/core	512KB/core	PCIe	300W	1,2 TFlop/s dupla pontosság
		OpenMP-nél csak 240		~100 cycles ~174 GB/s	~3 cycles	~10 cycles	~7 GB/s ~1 microsecond		2,4 TFlop/s egyszeres pontosság

Az Intel® 1. generációs Xeon Phi™ koprocesszorában található 61 processzor ugyanazt a 16GB memóriát éri el, de semmilyen módon se fér hozzá közvetlenül a hoszt gép fizikai memóriájához.
Budapest2 és Debrecen3-Phi gépeinkben számítási nodonként 2-2 Phi™ koprocesszor található.
A kártyákon speciális célú Linux kernel fut melyre a belépés limitáltan csak a rendszergazdák számára lehetséges SSH-n keresztül a hoszt gépről. Természetesen a kártyára belépve olyan érzésünk lehet, mintha egy sok processzoros gépen dolgoznánk, ami rendelkezik rengeteg memóriával. Ez így is van, de a hatékonyabb kihasználás érdekében az SSH-n keresztüli elérés nem lehetséges, csak a Slurm segítségével valamint a programozási nyelvek által támogatott csatornákon keresztül. Ennek részletezése lentebb megtalálható.

Használati módok

A koprocesszor használatai módjait két csoportba lehet sorolni. A natív és az offload használatot.
Ezen használati módokról a lenti fejezetekben több részletet is megtudhat. Az itt megfogalmazott "támogatás" csak a mi infrastruktúránkra érvényes!
A bemutatott példák c++ nyelven vannak, de természetesen c és fortran esetén is megtalálható az ekvivalens módszer.

a) nem támogatott mód

közvetlen natív használat

A natív mód arra szolgál, hogy a MIC kártyára lefordított alkalmazásunkat felmásoljuk a kártyára, majd belépünk SSH-n a kártyára és azon futtatjuk a bináris kódunkat.
Ez a legegyszerűbb használati mód, ugyanis nincs más dolgunk, csak a -mmic kapcsolót használni a fordításnál és a fordító egy máris a MIC architektúrára kész bináris készít nekünk.
Workflow a következő lenne:

1. Alkalmazás megírása. például hello.cc
2. Alkalmazás fordítása
   

   icpc -o hello-MIC -mmic hello.cc

3. Másolás a MIC kártyára
   

   scp hello-MIC mic0:

4. belépés a kártyára
   

   ssh mic0

5. Alkalmazás futtatása
   

   ./hello-MIC

b) támogatott, de nem preferált használat

automatikus kódátvitel

Ebben az esetben nem kell átmásolnunk a bináris alkalmazásunkat a kártyára sőt még csak be se kell lépnünk, mert ezt a két lépést egy beépített tool végzi el.
A tool eléréséhez szükséges parancs neve: micnativeloadex

Ebben az esetben a workflow a következőképpen módosul:

1. Alkalmazás megírása. például hello.cc
2. Alkalmazás fordítása
   

   icpc -o hello-MIC -mmic hello.cc

3. A tool futásához szükséges library-k megadása
   

   export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic

4. Alkalmazás futtatása
   

   micnativeloadex hello-MIC

BUDAPEST2[cn10] phi (0)$ export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic BUDAPEST2[cn10] phi (0)$ micnativeloadex hello-MIC Hello world! I have 244 logical cores. BUDAPEST2[cn10] phi (0)$

c) támogatott módok

Offload

Az Offload modell esetén szükséges a forráskód módosítása, mégpedig úgy, hogy azon kódrészletet, amit a koprocesszoron szeretnénk futtatni, ki kell egészíteni egy fordítási direktívával. Az alkalmazás fordításakor elegendő csak a CPU-ra való binárist készíteni mivel a modell végzi a megfelelő forrás részlet mozgatását és futtatását.

Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.

#include <stdio.h>
int main(int argc, char * argv[]) {
  printf("Hello world from host!\n");
#pragma offload target(mic)
  {
    printf("Hello world from coprocessor!\n"); fflush(stdout);
  }
  printf("Bye\n");
}

Ebben a példában az operációs rendszer választja ki, hogy melyik Phi kártyán fusson le a kódrészlet. Amennyiben mi szeretnénk explicite megmondani, hogy melyik kártyán fusson, akkor a fordítási direktívában a mic kulcsszó után kell tenni egy ":"-t és a kártya sorszámát [mi esetünkben 0 vagy 1]. Ekkor így nézne ki a direktíva: #pragma offload target(mic:1)

Amennyiben az offload szekcióban függvény hívás is szerepel, akkor szükségeltetik az függvény futásához szükséges változók adatterületének átmásolása a Phi kártyára, majd a lefutás után visszamásolni a hoszt gépre. Ezen másolást az offload_attribute push és pop paraméterei segítségével határozhatjuk meg.
A következő mátrix-vektor szorzás példánál az A mátrix globális változó, míg a b és c vektorok a függvény paraméterei.

//A változók felmásolásra kerülnek a Xeon Phi-re
#pragma offload_attribute(push, target(mic))
  double* A;
  void MatrixVector(double *b, double *c, double m, double n) {
    ...
    c[] = A[..]*b[]
  }
// felmásolt objektumok visszamásolása
#pragma offload_attribute(pop)

int main() {
  ...
#pragma offload target(mic) in(A[0:n*m])
// így is meg lehetne adni: #pragma offload target(mic:0) in(A : length(n*m))	        
  {
    MatrixVector(b, c, m, n);
  }
  ...
}

Ekkor minden memória területet 2x kell másolni, egyszer fel a kártyára, majd másodszor vissza a hoszt gépre, pedig elegendő lenne csak az eredmény vektort visszamásolni. Ezen optimalizálással egy későbbi fejezetben foglalkozunk. <p />

Futás során kaphatunk részletes riportot is, hogy pontosan milyen adatmozgatások történnek az alkalmazás futása során. Ehhez nincs szükség másra csak az OFFLOAD_REPORT környezeti változó megfelelő beállítása. A beállított értékek 1, 2 vagy 3 lehetnek, attól függően mennyire részletes riportot szeretnénk kapni.
beállítás:

export OFFLOAD_REPORT=3

példa kimenet a hello alkalmazásra:

Hello world from host!
[Offload] [MIC 0] [File]                    hello_offload.cc
[Offload] [MIC 0] [Line]                    5
[Offload] [MIC 0] [Tag]                     Tag 0
Hello world from coprocessor!
[Offload] [HOST]  [Tag 0] [CPU Time]        0.554653(seconds)
[Offload] [MIC 0] [Tag 0] [CPU->MIC Data]   0 (bytes)
[Offload] [MIC 0] [Tag 0] [MIC Time]        0.000140(seconds)
[Offload] [MIC 0] [Tag 0] [MIC->CPU Data]   0 (bytes)
Bye

Offload + OpenMP

Amennyiben nem csak 1 processzort szeretnénk használni a Phi kártyánkon és az alkalmazásunkban használt számítás megengedi a párhuzamosítást, akkor az OpenMP segítségével párhuzamosíthatjuk a megfelelő számítást.

Példánk maradjon ugyanúgy a mátrix-vektor szorzás. Amennyiben a teljes kártya adta kapacitást ki szeretnénk használni, akkor az OMP threadek számát 240-re kell beállítani. Azért nem 244-re, mivel 1 CPU-t dedikáltan csak a kommunikációra használ a rendszer és 1 CPU 4 hardver threaddel rendelkezik. Ekkor a példánk így változik:

#include <omp.h>

#pragma offload_attribute(push, target(mic))
void multiply(int n, int m, double* A, double* b, double* c){
  const int nCPUs = omp_get_max_threads();   //240
  omp_set_num_threads(nCPUs);
#pragma omp parallel for schedule(dynamic, 1)
  for ( int i = 0 ; i < m ; i++)
     for ( int j = 0 ; j < n ; j++)
       c[i] += A[i*n+j] * b[j];
}
#pragma offload_attribute(pop)

int main(){
  double * A = (double*) malloc(sizeof(double)*n*m);
  ...
#pragma offload target(mic) in(A:length(n*m))
     multiply(n, m, A, b, c);

Az alkalmazás fordításánál a fordítónak meg kell adni, hogy használhja az openmp library-ket is. Ezt a -qopenmp kapcsolóval tudjuk elérni.

icpc -o matrix-CPU -qopenmp matrix.cc

MPI + Offload + OpenMP

Abban az esetben, ha 1 számítási node nem elég, MPI segítségével össze tudunk kötni több nodet, heterogén rendszert alkotva, és azokon alkalmazható a fenti Offload + OpenMP modell.

Alkalmazások optimalizálása a kártyára

Ebben a fejezetben optimalizálási lehetőségeket mutatunk be, melyek segítségül szolgálhatnak az egyes alkalmazások hatékonyabb futása érdekében ezzel jobban kihasználni a rendelkezésre álló erőforrások adta számítási kapacitás.

Környezeti változók

Lehetséges csak a Phi kártyákon használt környezeti változók beállítása, amit az alkalmazásunk használni is tud.
Példánkban az OpenMP esetén használható szálak számát állíthatjuk be:

export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240

A beállított prefix tetszőleges karaktersorozat lehet, a lényeg, hogy ezt a beállított prefixet használjuk a környezeti változók esetén is!

Hosszredukció

Többnyire a fordítók, ha nem mondjuk meg explicite, a dupla pontosságú változatot választják a lehetséges változó- és függvény típusok közül. Amennyiben a számításunk nem igényli a dupla pontosságot, akkor használjuk a forrásban mindenhol az adott típus vagy függvény egyszeres pontosságú változatát, mellyel jelentős teljesítmény növekedést tudunk elérni. Továbbá az egyes műveletek futási ideje nem egyforma, így a műveletet helyettesíthetjük más művelettel, ami gyorsabb, akkor ismét javíthatunk a hatékonyságon.
Íme pár példa, amit alkalmazni lehet:

• "/" művelet helyett "*" -t használjunk, mert azt gyorsabban végre tudja hajtani a processzor
• pow => 1.0 / sqrt() , tetszőleges hatványozás, pl.: 3/2-ik hatvány
• hardver áltatl támogatott műveletek:
  • exp => exp2
  • log => log2
• dupla pontosság -> egyszeres pontosság lebegőpontos számok és műveletek esetén
  • 1.0 => 1.0f
  • pow => powf

Vektorizáció

Hatékony hossz választás

Olyan esetben jelent előnyt, ha például az n*n-es mátrix esetén n nem osztható 16-val, mert akkor a vektorizáció nem teljesen hajtható végre.

// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);

Ilyen esetben mindig marad egy kis vektor részlet, amire elemenként kell elvégezni a műveletet.
A megoldást az jelenti, hogy a memória allokálást kiterjesztjük és nagyobb memória területet foglalunk a változónak, így már alkalmazható lesz a vektorizáció, mivel a maradvány területen is el tudjuk végezni a műveletet egy lépésben.

int lda=n;
if (n % 16 != 0) lda += (16 - n%16); // most lda%16==0
float* A = _mm_malloc(sizeof(float)*n*lda, 64);

// használat
for (int i = 0; i < n; i++)
  for (int j = 0; j < n; j++)
    A[i*lda + j] = ...

"Héj ciklusok" kiküszöbölése

Ciklusok esetén lehetséges, hogy a vektorizációt nem teljes mértékben ismeri fel a fordító program. Ekkor segíthetünk a fordítónak egy fordítási direktívával (#pragma vector aligned), ha teljes mértékben biztosak vagyunk, hogy a ciklusunk vektorizálható!

#pragma vector aligned
  for (int j = 0; j < n; j++)
    A[i*lda + j] -= ...

Ettől jelentős futás idejű javulást érhetünk el, mivel ezzel a #pragma-val kikapcsoltuk a futás idejű ellenőrzést, de előfordulhatnak olyan esetek, főleg amikor nem jól mértük fel a ciklusunkat és nem sikerült a vektorizálás, hogy Segmentation fault hibára fogunk futni.

Regularizáció

Van egy ciklusunk, ami egy vektor elemeit módosítja. A processzor az adatokat "csokorba szedve" emeli be a belső kis méretű cache-be. Egyszerre természetesen több vektor elem is bekerül a a cache-be ezzel gyorsítva a feldolgozást. Ezen vektor iterációk valahogyan ráilleszkedik a vektorra.

Abban az esetben, amikor ez az illesztés nem teljesen tökéletes (lsd. piros elemek), keletkeznem héj ciklusok, melyeket a regularizációval is eliminálni lehet.
A feladat eredetileg valahogyan így nézhet ki:

for (int i = 0; i<n; i++)
  A[i]=...

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.

for (int b=0; b<n; b++) {
  const int jStart = b- b%16; // mindig osztható 16-val
  ...
#pragma simd
  for (int j = jStart; j<n; j++)
    A[j]=...