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.

5.0 KiB Raw Blame History