Elérkeztünk az utolsó részhez. Több gyakorlatot nem terveztem, de lesznek még CUDA-val kapcsolatos bejegyzések. Elég a bevezetőből, vágjunk bele.
A kódunk még messze nem tökéletes. Először csak rövid optimalizációkat végzünk. A szögfüggvényeknek például van a CUDA magokon futó változata. Ezt a math_functions.h fejléc állományban elfedik előlünk, hogy könnyű legyen a kód migrálása, de jobb, ha ezt magunk csináljuk, nem bízzuk egy fordítóra, amit nem mi írtunk :-) Ezért nézzük át a device_functions.h fejlécállományt. Találunk ott egy érdekes függvényt: __sincosf(). Tehát egy lépésben kiszámolhatjuk a színuszát és a koszinuszát ugyan annak a szögnek.
Általánosságban elmondható, hogy a két aláhúzással jelzett függvények eszköz kódok. Amikor csak tehetjük, használjuk őket. Szorzás helyett is nyugodtan használhatjuk az __fmul_rn függvényt. Ezek alapján az új distance függvényünk a következőképp fog kinézni:
__device__ float distance(float2 p1, float2 p2){
float dx = p1.x - p2.x;
float dy = p1.y - p2.y;
float x2 = __fmul_rn(dx,dx);
float y2 = __fmul_rn(dy,dy);
float su = __fadd_rn(x2,y2);
return( __fsqrt_rn(su) );
}
Megmondom őszintén, nem tudom, mit jelent a függvények végén az _rn. Van másik három rz, rd, ru is, nem tudom, hogy mi a különbség közöttük. Amint lesz róla információm, kijavítom ezt a bejegyzést.
__device__ void collision(float2 *p1, float2 *p2, float angle){
float p1x2 = __fmul_rn( p1[0].x, p1[0].x );
float p2x2 = __fmul_rn( p2[0].x, p2[0].x);
float p1y2 = __fmul_rn( p1[0].y, p1[0].y);
float p2y2 = __fmul_rn( p2[0].y, p2[0].y);
float m1 = __fsqrt_rn( p1x2 + p1y2);
float m2 = __fsqrt_rn( p2x2 + p2y2);
float d1 = atan2f( p1[0].y, p1[0].x);
float d2 = atan2f( p2[0].y, p2[0].x);
float sa, ca, sb, cb; /* sin and cos of d1 - angle */
__sincosf(d1 - angle, &sa, &ca);
__sincosf(d2 - angle, &sb, &cb);
float p1newx = __fmul_rn(m1,ca);
float p1newy = __fmul_rn(m1,sa);
float p2newx = __fmul_rn(m2,cb);
float p2newy = __fmul_rn(m2,sb);
__sincosf(angle + 3.14f / 2.0f, &sa, &ca);
__sincosf(angle, &sb, &cb);
p1[0].x = __fmul_rn(cb,p2newx) + __fmul_rn(ca, p1newy);
p1[0].y = __fmul_rn(sb,p2newx) + __fmul_rn(sa, p1newy);
p2[0].x = __fmul_rn(cb,p1newx) + __fmul_rn(ca, p2newy);
p2[0].y = __fmul_rn(sb,p1newx) + __fmul_rn(sa, p2newy);
}
A collision függvényünket is megváltoztathatjuk. Itt ráadásul áttekinthetőbbé tettük a paramétereket. A mutatók tömbként való használata nem véletlen, ha másképp használjuk őket, akkor "error: expression must have class type" hibaüzenetet kapunk fordítási időben.
Utolsóként a calcParticlePos függvényünk következik. A globális memória elérése nem túl ideális, egyrészt a lassabb elérés miatt, másrészt a konfliktusok miatt. Előfordulhat ugyanis, hogy több szál próbál meg módosítani egy memória területet.
A memória elérés gyorsítására lehet használni a megosztott memóriát. Annak bemutatására most nem vállalkozom.