# Párhuzamos primitívek

## 1. Hisztogram számítás

A feladathoz tartozó host kód a `primitives/Histogram.cpp` fájlban található.

### Globális memória

  Az implementációban csak a globális memória kerül felhasználásra, az alábbi pszeudokód szerint:

  ```
  id  := get_global_id(0)
  histogram[data[id]] := histogram[data[id]] + 1
  ```

  

 **1. Vizsgálja meg, hogy mi történik szinkronizáció nélkül! Helyes-e a kiszámolt eredmény?**

  Szinkronizáció nélkül a szálak egymásba akadnak végrehajtás során, így nem lesz helyes az eredmény.

 **2. Vizsgálja meg, hogy mi történik, ha a hisztogram megfelelő elemének növelése atomi módon
történik!**

  Ezzel a módszerrel egy elemhez csak egy szál férhet hozzá egyszerre, így jó értékeket fog számolni.

 **3. Vizsgálja meg, hogy hogyan változik a kernel futási ideje, ha a data tömb értékkészlete nagyon
szűk (pl. csak 0 és 1 értékek lehetségesek), valamint abban az esetben, ha megközelíti a data
tömb méretét!**

  A méréshez négy futás eredményeit használtam fel, az 5. feladat leírásában található segédosztály segítségével. Minden esetben 4096 elemű vektorokat elemez, az értékkészletük rendre 2, 128, 1024 és 4096 elemű. A vektorok elemei véletlenszerűen generáltak ezen értékkészleten belül. A következő ábrán a program kimenete szerepel, négysoronként csoportosítva a négy tesztesetre. Egy teszthez tartozó kimenet első során a futtatás paraméterei olvashatók (értékkészlet, bemeneti vektor méret), alatta a CPU, majd a GPU futtatási idő milliszekundumban, végül a teszt sikerességét jelző egy szavas visszajelzés a CPU és GPU eredményeinek összehasonlítása alapján.

  ```
  Histogram (type=gobal,data_size=4096,valueSet=2)
      CPU [time] 0.23806
      GPU [time] 0.009472 ms
        Success
  Histogram (type=gobal,data_size=4096,valueSet=128)
      CPU [time] 0.23646
      GPU [time] 0.005728 ms
          Success
  Histogram (type=gobal,data_size=4096,valueSet=1024)
      CPU [time] 0.22818
      GPU [time] 0.00416
          Success
  Histogram (type=gobal,data_size=4096,valueSet=4096)
      CPU [time] 0.22962
      GPU [time] 0.003968 ms
          Success
  ```

  A GPU sorok összehasonlításával leolvasható, hogy a kisebb értékkészlet jelentősen hosszabb futási időt erdményezett. Ez visszavezethető arra, hogy a szálaknak várakozniuk kellett egymásra, amíg a hisztogram értékeit frissítették.

### Lokális memória

Az implementációban a munkacsoportok először lokális memóriába írják számításaikat, majd azzal frissítik a globális adatokat. Az algoritmust alkotó pszeudokód;

```
ID  := get_global_id(0)
LID := get_local_id(0)

IF LID < histogramSize DO:
  lhistogram[LID] := 0
BARRIER

lhistogram[data[id]] := lhistogram[data[id]] + 1

BARRIER

IF LID < histogramSize DO:
  histogram[LID] = lhistogram[LID]
```

**1. Vizsgálja meg, hogy mi történik barrier hívások használata nélkül. Helyes-e az eredmény?**

Barrier hívások nélkül a kernel helytelen eredményt ad.

 **2. Hasonlítsa össze teljesítmény szempontjából a globális és a lokális memóriában végzett
hisztogram számtást!**

  Az összehasonlításhoz egységesen 4096 elemű, de változó értékkészletű vektorok hisztogramját számítottam ki. Az alábbi táblázatban az oszlopok az értékkészlet nagyságát jelölik, a sorok a globális, illetve a lokáis kerneleket, a cellák pedig az egyes futási időket egy adott típusó kernelhez egy értékkészleten.

  |    **/**   |    **2**    |   **128**   |   **200**   |   **250**   |
  |    ---     |     ---     |     ---     |     ---     |     ---     |
  | **local**  | 0.006592 ms | 0.004064 ms | 0.004512 ms | 0.004448 ms |
  | **global** | 0.00784 ms  | 0.0064 ms   | 0.0064 ms   | 0.005952 ms |

A futási időkből látszik, hogy a lokális memóriát használó kernel minden értékkészletre gyorsabb, viszont a lokális és globális kernelek is gyorsabban futnak, ha ez az értékkészlet nagyobb.

## 2. Összegzés redukcióval

A feladathoz tartozó host kód a `primitives/ReduceAdd.cpp` fájlban található. Az implementáció során a programs.cl forrásfájl pszeudokódja alapján dolgoztam, és feltételeztem, hogy a tömb mérete legfeljebb a workgroup megengedett maximális mérete.

## 3. Exkluzív prefix összeg scan operátorral

A feladathoz tartozó host kód a `primitives/ExclusiveScan.cpp` fájlban található. Az implementáció során a programs.cl forrásfájl pszeudokódja alapján dolgoztam, kicsit módosítva azt egy plusz barrier hozzáadásával a FOR ciklusban;

```
ID := get_global_id(0)
VAL := 0
IF ID > 0 THEN VAL = data[ID - 1]
          ELSE VAL = 0
BARRIER
data[ID] = VAL

FOR s = 1; s < get_global_size(0); s *= 2 DO:
  tmp := data[ID]
  BARRIER
  IF ( ID + s < get_global_size(0) THEN
    data[ID + s] += tmp;
  BARRIER

IF(ID = 0) THEN data[ID] = 0;
```

A 2. feladathoz hasonlóan itt is feltételeztem, hogy a tömb mérete legfeljebb a workgroup megengedett maximális mérete.

## 4.  Elemek számlálása compact operátorral

A feladathoz tartozó host kód a `primitives/Compact.cpp` fájlban található. A Compact operátort a feladat specifikációja szerint három kernel összefűzésével készítettem el;

 * Map - A bementi vektor értékeihez nullát, vagy egyet rendel annak függvényében, hogy az kisebb-e 50-nél.
 * Exkluzív scan - Az előbbi lépés kimenetét felhasználva kiszámítja, melyik indexre fognak kerülni a nem ritka eredménytömbben az 50-nél kisebb elemek.
 * Scatter - Előző két lépés eredményeit felhasználva előállítja a nem ritka eredmény tömböt. 

Mivel az eredménytömb kisebb is lehet, mint a bemeneti, akkor ha az eddigiekhez hasonlóan a bemeneti tömb méretét olvasom vissza az OpenCL bufferből, felesleges adatokat is mozgatok. Ennek elkerülése érdekében végeztem pár módosítást a mellékelt pszeudokód által leírt algoritmuson. Az exkluzív scan lépésben a globálisan nulladik azonosítójú szál a tömb legutolsó elemével felülírja a saját értékét, így a tömb első eleme az 50-nél kisebb értékű elemek számát fogja jelezni. A későbbiekben a nulladik elem redundáns ebben a tömbben, hiszen ha a predikátum 0, nem kerül bele az eredménybe, ha viszont 1, garantáltam a 0. helyre fog bekerülni. Ennek megfelelően a scatter lépést is módosítottam, hogy ne legyen szükség a scan nulladik elemére;

```
ID := get_global_id(0)
VALUE := data[ID]
BARRIER
IF pred[ID] == 1 THEN
  IF ID == 0 THEN
    data[0] = VALUE
  ELSE
    data[prefSum[ID]] = VALUE
```

## 5. Időmérés

Megoldás során a feladatkiírásban megadott kódot használtam. A host kód mérésére szolgáló Timer osztály a Common.cpp állományba került, az OpenCL device kód futási idejének mérésére szolgáló segédfüggvény pedig az OpenCLHandler osztály egyik függvénye lett. A mérés során az alábbi kerneleket, illetve a hozzájuk tartozó host kódot futtattam;

 * Square - 4096 elemű vektor
 * Hisztogram 1 - globális, 4096 elemű vektor, 32 elemű értékkészlet
 * Hisztogram 2 - lokális, 4096 elemű vektor, 32 elemű értékkészlet
 * Redukció - 1024 elemű vektor
 * Exkluzív scan - 1024 elemű vektor
 * Compact - 1024 elemű vektor

Az alábbi táblázat a mérési eredményeket szemlélteti;

|     kernel    | CPU (ms) | GPU (ms) |
|  -----------  |  ------  |  ------  |
| Square        | 1.2394   | 0.006304 |
| Hisztogram 1  | 0.23674  | 0.006432 |
| Hisztogram 2  | 0.2357   | 0.004096 |
| Redukció      | 0.01558  | 0.00784  |
| Exkluzív scan | 8.0889   | 0.013184 |
| Compact       | 0.17338  | 0.00432  |

Az eredményekből is látszik, hogy valóban nagy méretű sebességnövekedés érhető el masszív párhuzamosítással.

## TODO
  * Elvégezni a mérést _nem_ debug módban