HTML

Az élet kódjai

Csináld maga. Senki nem csinálja meg helyetted.

Friss topikok

CUDA (4. rész)

2011.02.04. 12:27 Travis.CG

Ő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.

Szólj hozzá!

Címkék: programozás

A bejegyzés trackback címe:

https://cybernetic.blog.hu/api/trackback/id/tr512638367

Kommentek:

A hozzászólások a vonatkozó jogszabályok  értelmében felhasználói tartalomnak minősülnek, értük a szolgáltatás technikai  üzemeltetője semmilyen felelősséget nem vállal, azokat nem ellenőrzi. Kifogás esetén forduljon a blog szerkesztőjéhez. Részletek a  Felhasználási feltételekben és az adatvédelmi tájékoztatóban.

Nincsenek hozzászólások.
süti beállítások módosítása