09/12/03 07:38:09
あ、失礼 coalescedになってなかった。こうかな??
//4回CUDA_Funcを呼び出す方向で。
for (int n = 0; n < 1024 / 256; n ++)
{
CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, n);
}
//nは「縦」の分割数
//512スレッドが連続した512バイトを取り込む。二回動くと、1ピクセル×横に1024ピクセルを処理。
//上に向かって縦256回回る（k)
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int n)
{
for (k=0; k < 256; k++)
{
//動き出す起点は各スレッドで1バイトずつずれてる。
int address = n*1024*256 + k*1024 + threadIDx.x;
//1024バイトを512スレッドで処理するので、二回。
pDest[address] = ((int)pSrc1[address]+(int)pSrc2[address])>>1;
pDest[address+512] = ((int)pSrc1[address+512]+(int)pSrc2[address+512])>>1;
}
}
実際書いて動かさないと良く分からないすな。グレーのビットマップ二つ用意してやってみる形かな。