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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===
A vektorizáció azon fordítási időben eldönthető optimalizációs faktor, amely nagy mértékben függ attól a hardver felépítésétől, amire éppen a fordítás történni fog. Legfontosabb befolyásoló tényezők a processzor cache méretei illetve a vektor műveletekért felelős aritmetikai egységének jellemzői mivel ezeket használja fel a fordító az optimalizálási folyamata során. 
Az automatikus vektorizáció minimum -O2 szinttől érvényes, de mivel általában ez a szint az alapértelmezett, ezért minden fordításnál alkalmazásra kerül. 
A fordító programnak a következő kapcsolókat lehet megadni, melyek segítenek a minél jobb vektorizációban:
* ''-qopt-report=<szám>'' ; az optimalizációs jelentés részletességét lehet megadni 0-tól 5-ig, ahol az 5-s szint a legrészletesebbet jelenti
* ''-qopt-report-phase:vec'' ; a vektorizációról szeretnék jelentést kapni
* ''-opt-report-file=<fájlnév>'' ; megadható a riportfájl neve

<pre>
icpc autovec.cc -mmic -qopt-report=5 -opt-report-file=vectorization.oMIC.optrpt -qopt-report-phase:vec
</pre>

====SIMD====
Az adatok hatékonyabb feldolgozására az egyik megvalósítás a SIMD (Single Instruction Multiple Data), amikor is a különálló adatokat úgy kezeli, mintha az vektorok volnának és azon végzi el a vektor-műveletet. Ezzel egy a hardver által támogatott műveletet végzünk egy lépésben egy nagyobb adatmennyiségen, ezzel párhuzamosítva a feldolgozási folyamatot. Ezzel a módszerrel nagyon jó adat szintű párhuzamosítást lehet elérni. 

[[Fájl:Simd.png|SIMD]]

Főbb támogatott műveletek:
* aritmetikai műveletek, bináris operátorok
* összehasonlítások
* konverziók és típus kasztolások
* bit-műveletek: NOT, AND, OR, XOR, XAND

Használata: '''simd''' fordítási direktívával lehetséges 
Például:
* egyszerű vektorizáció
<pre>
#pragma simd
for (int i = 0; i < n; i++)
A[i] += B[i];
</pre>

* adat párhuzamosítás OpenMP segítségével
<pre>
#pragma omp parallel for simd
for (int i = 0; i < n; i++)
DoSomWork(A[i]);
</pre>

====AoS -> SoA====
A kiindulás probléma az, hogy a struktúrákat tartalmazó vektort nem lehet jól vektorizálni, mivel a cache-ben tárolt adatok nagy többsége feleslegesen van bent, mivel azon adatokat csak egy későbbi ciklus futás során fogja az alkalmazás feldolgozni. Tehát ezen adatok jelenleg feleslegesen foglalják a cache területet, ahelyett, hogy hasznos adatot tárolnánk ott.
<pre>
struct MyStructure {
float x, y, z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
for (int i = 0; i < nItem; i++) {
...
for (int j = 0; j < nItem; j++) {
const float dx = Items[j].x - Items[i].x;
const float dy = Items[j].y - Items[i].y;
const float dz = Items[j].z - Items[i].z;
...
}
}
...
}
</pre>

A vektorban a struktúrák tárolása helyett minden elemnek egy dedikált vektort allokálunk és így könnyen tudja a fordító vektorizálni azt. Természetesen segítségként a vektor feldolgozását "tömbösítettük" a tileSize méret megadásával illetve a ''unroll'' fordítási direktíva segítségével.
<pre>
struct MyStructureSet {
float *x, *y, *z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
const int tileSize = 16;
#pragma omp parallel for schedule(guided)
for (int ii = 0; ii < nItem; ii+= tileSize) {
...
#pragma unroll(tileSize)
for (int j = 0; j < nItem; j++) {
#pragma unroll(tileSize)
for (int i = ii; i < ii + tileSize; i++) {
const float dx = Items.x[j] - Items.x[i];
const float dy = Items.y[j] - Items.y[i];
const float dz = Items.z[j] - Items.z[i];
...
}
}
}
...
}
</pre>

====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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.

[[Fájl:regularization.png|Vektor regularizáció]]

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:
<pre>
for (int i = 0; i<n; i++)
A[i]=...
</pre>

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.
<pre>
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]=...

</pre>

===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Fájl:Simd.png

2016-05-05T07:22:36Z

Szigeti(AT)niif.hu:

Intel Xeon Phi

2016-05-04T11:17:55Z

Szigeti(AT)niif.hu: /* Vektorizáció */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===
A vektorizáció azon fordítási időben eldönthető optimalizációs faktor, amely nagy mértékben függ attól a hardver felépítésétől, amire éppen a fordítás történni fog. Legfontosabb befolyásoló tényezők a processzor cache méretei illetve a vektor műveletekért felelős aritmetikai egységének jellemzői mivel ezeket használja fel a fordító az optimalizálási folyamata során. 
Az automatikus vektorizáció minimum -O2 szinttől érvényes, de mivel általában ez a szint az alapértelmezett, ezért minden fordításnál alkalmazásra kerül. 
A fordító programnak a következő kapcsolókat lehet megadni, melyek segítenek a minél jobb vektorizációban:
* ''-qopt-report=<szám>'' ; az optimalizációs jelentés részletességét lehet megadni 0-tól 5-ig, ahol az 5-s szint a legrészletesebbet jelenti
* ''-qopt-report-phase:vec'' ; a vektorizációról szeretnék jelentést kapni
* ''-opt-report-file=<fájlnév>'' ; megadható a riportfájl neve

<pre>
icpc autovec.cc -mmic -qopt-report=5 -opt-report-file=vectorization.oMIC.optrpt -qopt-report-phase:vec
</pre>

====SIMD====
Az adatok hatékonyabb feldolgozására az egyik megvalósítás a SIMD (Single Instruction Multiple Data), amikor is a különálló adatokat úgy kezeli, mintha az vektorok volnának és azon végzi el a vektor-műveletet. Ezzel egy a hardver által támogatott műveletet végzünk egy lépésben egy nagyobb adatmennyiségen, ezzel párhuzamosítva a feldolgozási folyamatot. Ezzel a módszerrel nagyon jó adat szintű párhuzamosítást lehet elérni. 
Főbb támogatott műveletek:
* aritmetikai műveletek, bináris operátorok
* összehasonlítások
* konverziók és típus kasztolások
* bit-műveletek: NOT, AND, OR, XOR, XAND

Használata: '''simd''' fordítási direktívával lehetséges 
Például:
* egyszerű vektorizáció
<pre>
#pragma simd
for (int i = 0; i < n; i++)
A[i] += B[i];
</pre>

* adat párhuzamosítás OpenMP segítségével
<pre>
#pragma omp parallel for simd
for (int i = 0; i < n; i++)
DoSomWork(A[i]);
</pre>

====AoS -> SoA====
A kiindulás probléma az, hogy a struktúrákat tartalmazó vektort nem lehet jól vektorizálni, mivel a cache-ben tárolt adatok nagy többsége feleslegesen van bent, mivel azon adatokat csak egy későbbi ciklus futás során fogja az alkalmazás feldolgozni. Tehát ezen adatok jelenleg feleslegesen foglalják a cache területet, ahelyett, hogy hasznos adatot tárolnánk ott.
<pre>
struct MyStructure {
float x, y, z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
for (int i = 0; i < nItem; i++) {
...
for (int j = 0; j < nItem; j++) {
const float dx = Items[j].x - Items[i].x;
const float dy = Items[j].y - Items[i].y;
const float dz = Items[j].z - Items[i].z;
...
}
}
...
}
</pre>

A vektorban a struktúrák tárolása helyett minden elemnek egy dedikált vektort allokálunk és így könnyen tudja a fordító vektorizálni azt. Természetesen segítségként a vektor feldolgozását "tömbösítettük" a tileSize méret megadásával illetve a ''unroll'' fordítási direktíva segítségével.
<pre>
struct MyStructureSet {
float *x, *y, *z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
const int tileSize = 16;
#pragma omp parallel for schedule(guided)
for (int ii = 0; ii < nItem; ii+= tileSize) {
...
#pragma unroll(tileSize)
for (int j = 0; j < nItem; j++) {
#pragma unroll(tileSize)
for (int i = ii; i < ii + tileSize; i++) {
const float dx = Items.x[j] - Items.x[i];
const float dy = Items.y[j] - Items.y[i];
const float dz = Items.z[j] - Items.z[i];
...
}
}
}
...
}
</pre>

====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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.

[[Fájl:regularization.png|Vektor regularizáció]]

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:
<pre>
for (int i = 0; i<n; i++)
A[i]=...
</pre>

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.
<pre>
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]=...

</pre>

===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-04T09:05:22Z

Szigeti(AT)niif.hu: /* SIMD */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
Az adatok hatékonyabb feldolgozására az egyik megvalósítás a SIMD (Single Instruction Multiple Data), amikor is a különálló adatokat úgy kezeli, mintha az vektorok volnának és azon végzi el a vektor-műveletet. Ezzel egy a hardver által támogatott műveletet végzünk egy lépésben egy nagyobb adatmennyiségen, ezzel párhuzamosítva a feldolgozási folyamatot. Ezzel a módszerrel nagyon jó adat szintű párhuzamosítást lehet elérni. 
Főbb támogatott műveletek:
* aritmetikai műveletek, bináris operátorok
* összehasonlítások
* konverziók és típus kasztolások
* bit-műveletek: NOT, AND, OR, XOR, XAND

Használata: '''simd''' fordítási direktívával lehetséges 
Például:
* egyszerű vektorizáció
<pre>
#pragma simd
for (int i = 0; i < n; i++)
A[i] += B[i];
</pre>

* adat párhuzamosítás OpenMP segítségével
<pre>
#pragma omp parallel for simd
for (int i = 0; i < n; i++)
DoSomWork(A[i]);
</pre>

====AoS -> SoA====
A kiindulás probléma az, hogy a struktúrákat tartalmazó vektort nem lehet jól vektorizálni, mivel a cache-ben tárolt adatok nagy többsége feleslegesen van bent, mivel azon adatokat csak egy későbbi ciklus futás során fogja az alkalmazás feldolgozni. Tehát ezen adatok jelenleg feleslegesen foglalják a cache területet, ahelyett, hogy hasznos adatot tárolnánk ott.
<pre>
struct MyStructure {
float x, y, z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
for (int i = 0; i < nItem; i++) {
...
for (int j = 0; j < nItem; j++) {
const float dx = Items[j].x - Items[i].x;
const float dy = Items[j].y - Items[i].y;
const float dz = Items[j].z - Items[i].z;
...
}
}
...
}
</pre>

A vektorban a struktúrák tárolása helyett minden elemnek egy dedikált vektort allokálunk és így könnyen tudja a fordító vektorizálni azt. Természetesen segítségként a vektor feldolgozását "tömbösítettük" a tileSize méret megadásával illetve a ''unroll'' fordítási direktíva segítségével.
<pre>
struct MyStructureSet {
float *x, *y, *z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
const int tileSize = 16;
#pragma omp parallel for schedule(guided)
for (int ii = 0; ii < nItem; ii+= tileSize) {
...
#pragma unroll(tileSize)
for (int j = 0; j < nItem; j++) {
#pragma unroll(tileSize)
for (int i = ii; i < ii + tileSize; i++) {
const float dx = Items.x[j] - Items.x[i];
const float dy = Items.y[j] - Items.y[i];
const float dz = Items.z[j] - Items.z[i];
...
}
}
}
...
}
</pre>

====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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.

[[Fájl:regularization.png|Vektor regularizáció]]

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:
<pre>
for (int i = 0; i<n; i++)
A[i]=...
</pre>

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.
<pre>
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]=...

</pre>

===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-04T07:58:54Z

Szigeti(AT)niif.hu: /* AoS -> SoA */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
A kiindulás probléma az, hogy a struktúrákat tartalmazó vektort nem lehet jól vektorizálni, mivel a cache-ben tárolt adatok nagy többsége feleslegesen van bent, mivel azon adatokat csak egy későbbi ciklus futás során fogja az alkalmazás feldolgozni. Tehát ezen adatok jelenleg feleslegesen foglalják a cache területet, ahelyett, hogy hasznos adatot tárolnánk ott.
<pre>
struct MyStructure {
float x, y, z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
for (int i = 0; i < nItem; i++) {
...
for (int j = 0; j < nItem; j++) {
const float dx = Items[j].x - Items[i].x;
const float dy = Items[j].y - Items[i].y;
const float dz = Items[j].z - Items[i].z;
...
}
}
...
}
</pre>

A vektorban a struktúrák tárolása helyett minden elemnek egy dedikált vektort allokálunk és így könnyen tudja a fordító vektorizálni azt. Természetesen segítségként a vektor feldolgozását "tömbösítettük" a tileSize méret megadásával illetve a ''unroll'' fordítási direktíva segítségével.
<pre>
struct MyStructureSet {
float *x, *y, *z;
}

void My_Function(cons int nItem, MyStructure* const Items, const float delta) {
const int tileSize = 16;
#pragma omp parallel for schedule(guided)
for (int ii = 0; ii < nItem; ii+= tileSize) {
...
#pragma unroll(tileSize)
for (int j = 0; j < nItem; j++) {
#pragma unroll(tileSize)
for (int i = ii; i < ii + tileSize; i++) {
const float dx = Items.x[j] - Items.x[i];
const float dy = Items.y[j] - Items.y[i];
const float dz = Items.z[j] - Items.z[i];
...
}
}
}
...
}
</pre>

====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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.

[[Fájl:regularization.png|Vektor regularizáció]]

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:
<pre>
for (int i = 0; i<n; i++)
A[i]=...
</pre>

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.
<pre>
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]=...

</pre>

===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-03T12:42:27Z

Szigeti(AT)niif.hu: /* Regularizáció */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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.

[[Fájl:regularization.png|Vektor regularizáció]]

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:
<pre>
for (int i = 0; i<n; i++)
A[i]=...
</pre>

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.
<pre>
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]=...

</pre>

===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Fájl:Regularization.png

2016-05-03T12:41:26Z

Szigeti(AT)niif.hu:

Intel Xeon Phi

2016-05-03T12:17:07Z

Szigeti(AT)niif.hu: /* Regularizáció */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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.

[[Fájl:regularozation.png|Vektor regularizáció]]

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:
<pre>
for (int i = 0; i<n; i++)
A[i]=...
</pre>

Ez a következő átalakítással már alkalmas vektorizációra és nem lesz "peel loop" probléma.
<pre>
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]=...

</pre>

===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-03T12:00:39Z

Szigeti(AT)niif.hu: /* "Héj ciklusok" kiküszöbölése */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="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ó!
<pre>
#pragma vector aligned
for (int j = 0; j < n; j++)
A[i*lda + j] -= ...
</pre>
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ó====
===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-03T11:50:54Z

Szigeti(AT)niif.hu: /* Hatékony hossz választás */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
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] = ...
</pre>

===="Héj ciklusok" kiküszöbölése====
Ciklusok esetén lehetséges, hogy a vektorizáció során

====Regularizáció====
===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-03T11:49:48Z

Szigeti(AT)niif.hu: /* "Héj ciklusok" kiköszöbölése */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
int lda=n;
if (n % 16 != 0) lda += (16 - n%16); // most lda%16==0
float* A = _mm_malloc(sizeof(float)*n*lda, 64);
</pre>

===="Héj ciklusok" kiküszöbölése====
Ciklusok esetén lehetséges, hogy a vektorizáció során

====Regularizáció====
===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====

Intel Xeon Phi

2016-05-03T11:40:09Z

Szigeti(AT)niif.hu: /* Hatékony hossz választás */

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
<gallery>
Fájl:Xeon-phi-family-rwd.png.rendition.intel.web.720.405.png|
</gallery>

{| class="wikitable" border="1"
|- align="center"
! CPU
! threads/Core
!max threads
! CPU Freq
! Memória
! L1 cache
! L2 cache
! Interfész
! Ptot
! Számítási kapacitás
|- align="center"
| 61
| 4 Hardver
| 244
| 1238 GHz
| 16GB GDDR5
| 32KB/core
| 512KB/core
| PCIe
| 300W
| 1,2 TFlop/s dupla pontosság
|- align="center"
|
|
| 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_klaszter|Budapest2]] és [[HPC#Szupersz.C3.A1m.C3.ADt.C3.B3g.C3.A9peink_.C3.B6sszehasonl.C3.ADt.C3.A1sa|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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# Másolás a MIC kártyára <pre>scp hello-MIC mic0:</pre>
# belépés a kártyára <pre>ssh mic0</pre>
# Alkalmazás futtatása <pre>./hello-MIC</pre>

===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:
# Alkalmazás megírása. például [[Média:hello.cc|''hello.cc'']]
# Alkalmazás fordítása <pre>icpc -o hello-MIC -mmic hello.cc</pre>
# A tool futásához szükséges library-k megadása <pre>export SINK_LD_LIBRARY_PATH=/opt/intel/composerxe/compiler/lib/mic</pre>
# Alkalmazás futtatása <pre>micnativeloadex hello-MIC</pre>

:{| class="wikitable"
|-
| 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. 
[[Fájl:offload.png|Offload modell]] 
Példánkban jól látható, hogy a második kiíratás már a koprocesszoron fut.
<pre>
#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");
}
</pre>

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.
<pre>
//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);
}
...
}
</pre>

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. 

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:
<pre>
export OFFLOAD_REPORT=3
</pre>
példa kimenet a hello alkalmazásra:
<pre>
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
</pre>

====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. 
[[Fájl:Offload+openmp.png|Offload modell és OpenMP]] 
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:
<pre>
#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);
</pre>

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.
<pre>
icpc -o matrix-CPU -qopenmp matrix.cc
</pre>

====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. 
[[Fájl:Mpi+offload+openmp.png|MPI elosztás valamint Offload modell és OpenMP]] 

==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:
<pre>
export MIC_ENV_PREFIX=XEONPHI
export OMP_NUM_THREADS=12
export XEONPHI_OMP_NUM_THREADS=240
</pre>
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ó===

====SIMD====
====AoS -> SoA====
====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.
<pre>
// n%16 != 0
float* A = _mm_malloc(sizeof(float)*n*n, 64);
</pre>
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.
<pre>
int lda=n;
if (n % 16 != 0) lda += (16 - n%16); // most lda%16==0
float* A = _mm_malloc(sizeof(float)*n*lda, 64);
</pre>

===="Héj ciklusok" kiköszöbölése====
====Regularizáció====
===Párhuzamosítás===
====Párhuzamos szálak ütemezése====
====Ciklus kifejtés====
====Automatikus ciklusösszevonás====