Őszintén szólva, az előző programunkban több hiba is van, de talán segít megérteni, hogy a CUDA programozás új gondolkodást kíván. Nem elég, hogy a ciklusokat kernelekre cseréljük. Az ütközős program a struktúrális programozás szellemében készült. Nem csak programhiba, hanem metodikai hiba is van benne. Most nézzük, hogyan tudjuk kijavítani ezeket.
Az SDK tartalmaz egy hibakereső programot cuda-gdb néven. A --device-debug opció segítségével hibakereső információk kerülnek a fordított kódba, szintje 0-3-ig terjedhet. Ezután használhatjuk a cuda-gdb programot, ami szintén az SDK része. Sajnos grafikus alkalmazásoknál nem használható, így használatától el kell tekintenünk.
A második, diagnosztikai módszer, amit minden bizonnyal mindenki ismert, a megjegyzések elhintése vagy a képernyőre vagy egy fájlba. A CUDA SDK tartalmaz egy simplePrintf példaprogramot, amivel üzeneteket helyezhetünk a képernyőre. Személy szerint nem kedvelem, de ettől még lehet, hogy jó hasznát veszi az ember.
Kreatív módszerekkel viszont élhetünk. Még annak idején, amikor nem voltak ilyen nagyszerű shader debugger programok, a programozók a beépített OpenGL változóknak adtak feltűnő értékeket, ha valami gyanús történt. Pl. fragment shader esetén vörös színt, vertex shader esetén a csúcspontot eltolták jobbra. Hasonló módszert mi is alkalmazhatunk. Adjunk hozzá még egy cudaGraphicsResource-t, amiben mondjuk a színeket tároljuk. Ha valami furcsaság történik, változtassuk meg a részecske színét!
A teljes programsort most nem írom ide, majd ha mégis igény lesz rá, akkor belinkelem. A kernelt viszont bemásolom. A lényegi műveleteket úgyis az végzi. A c tömbben a színeket tárolom:
__global__ void calcParticlePos(float2 *p, float2 *v, float3 *c){
int index = threadIdx.x;
int i;
/* Collision with the wall */
if(p[index].x > 1.0f){
p[index].x = 1.0f;
v[index].x *= -1.0f;
}
if(p[index].x < -1.0f){
p[index].x = -1.0f;
v[index].x *= -1.0f;
}
if(p[index].y > 1.0f){
p[index].y = 1.0f;
v[index].y *= -1.0f;
}
if(p[index].y < -1.0f){
p[index].y = -1.0f;
v[index].y *= -1.0f;
}
/* Collision with each other */
for(i = 0; i < PARTICLENUM; i++){
if(i == index) continue;
if(distance(p[i].x, p[i].y, p[index].x, p[index].y) < COLLIDEDIST){
c[i].x = 1.0f;
c[i].y += 0.01f;
c[i].z += 0.01f;
c[index].x = 1.0f;
c[index].y += 0.01f;
c[index].z += 0.01f;
float angle = atan2f(p[i].y - p[index].y, p[i].x - p[index].x);
collision(&v[i].x, &v[i].y, &v[index].x, &v[index].y, angle);
}
}
p[index].x = p[index].x + v[index].x;
p[index].y = p[index].y + v[index].y;
}
A program futása után már érezhető a probléma. A részecskék a kelleténél többször ütköznek. Ahelyett, hogy a részecskék 100 ütközés után érnék el a fehér színt, már az első ütközés alkalmával kifehérednek. Ennek az oka, hogy a CUDA magok igaz párhuzamosan futnak, de nem azonos idő alatt fejezik be a futást. A másik ok, amit szemfülesebb olvasók már bizonyára észrevettek, hogy a ciklusunk rossz! Minden egyes ütközés kétszer zajlik le. Tegyük fel, hogy a 231. részecske és a 42. részecske ütközését vizsgáljuk. Először lefut a 42. CUDA kernel, ahol a ciklus segítségével megtalálja az 231. részecskét, amivel ütközött. De a 231. CUDA kernel is lefut, és ő is megtalálja a 42. részecskét, amivel ütközik. A vizsgálat kétszer lesz igaz. Írjuk át a for ciklust a következőkre: for(i = index + 1; i < PARTICLENUM; i++). Ezzel az utána következő sorban található feltételt is kiváltottuk.
Ha lefuttatjuk a módosított kódot, akkor egy másik anomáliával találkozhatunk: egyes részecskék "összeragadnak" és egymás körül keringenek. Ez azért van, mert az ütközés után nem tudnak elég távol kerülni egymástól, ezért a következő ciklusban is ütköznek.
Mint említettem, az egyes szálak nem azonos sebességgel futnak. A fenti kódunkban ez okozhat problémákat. Kiküszöbölésükre találták ki a __syncthreads() függvényt. Használata csökkenti a teljesítményt, de segítségével elérhetjük, hogy a szálak nem versenyezzenek. Közös memóriaterület használata esetén nagyon hasznos tud lenni. Helyezzünk el egy ilyen függvényt nyugodtan a ciklus előtt.
A következő részben optimalizálni fogjuk a programunkat.
