Az új hatékonyabb kódunkhoz gyakorlatilag csak a kernelt kell módosítanunk.
Vegyük észre ugyanis, hogy ha a globális memóriában végzünk atomikus műveleteket az nem túl hatékony.
Egyrészt, ha feltesszük, hogy a számítógépben van 32 multiprocesszor, amiken amik egyszerre 32-32 aktív szálat kezelnek, akkor 32*32=1024 szál próbál növelni 256 hisztogrmam számlálót. Ez legjobb esetben is számlálónként egy négyszeres konfliktus, aminek feloldása 4*1200 = 4800 órajelt vesz igénybe.
Van viszont a multiprocesszorokban egy közös memória is, amiben tudunk tömböket tárolni, és blokkon belül közösen kezelni. Így tehát, ha a blokkok szálai a közös memóriában előállítanának egy részeredményt, akkor egyszerre csak 32 szál között lehetne konfliktus, amik 256 értéket kezelnek. Itt mégy figyelembe kell venni, hogy a közös memória 32 bankba van szervezve tehát 32 szál növel 32 bankot, de az arány akkor is 32-32 lesz.
Ehhez jön, hogy a közös memória elérése sokkal gyorsabb, mint a globális memóriáé, ugyanis a közös memória bankjai 4 órajel alatt elérhetőek (szemben a globális memória 600 órajeles késleltetésével). A számítás tehát kihozza, hogy statisztikailag egy-egy növeléshes 8 órajel lég lehet a korábbi 4800 helyett.
Az új kernel tehát először is közös memóriában tárol egy rész-hisztogramot.
__shared__ unsigned int temp[256];
temp[threadIdx.x] = 0;
__syncthreads();
A közös hisztogram foglalása és nullázása után kiszámítjuk a részeredményt.
int i = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
while (i < size) {
atomicAdd( &temp[buffer[i]], 1 );
i += stride;
}
Végül pedig csak annyi feladatunk marad, hogy az számokat a közös memóriából kiírjuk a globális memóriába.
__syncthreads();
atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );
Ez már nem annyira időigényes feladat, mivel hisztogram számláló sokkal kevesebb van, mint pixel.
A számításban most kihasználtuk, hogy pont annyi szál fut, ahány hisztogram elem van, így egy-egy érték kiírása szálanként pont megfelelő.