- 238 名前：デフォルトの名無しさん mailto:sage [2009/12/03(木) 06:59:13 ]
- >>229
ごめん元のソースが何したいのかよく分からないんだけど、こんなことしたいわけ？
1024ピクセル×1024ピクセルの二つの領域の明るさ平均を取るみたいな？
CPUなら、縦座標が外側ループで1024回×横座標が内側ループで1024回、回ると。
//device側で1MB×3(dest, src1, src2)をアロケートして、hostからsrc1,src2にデータをコピー
//512スレッドを起動するなら、外側ループの回る回数はCPUだと1024、GPUだと1024÷スレッド数512
for (int n = 0; n < 1024 / 512; n ++)
{
CUDA_Func<<<1, 512, 0>>>(dest, src1, src2, nPos);
}
//deviceからHostへdestをコピー
=====
__global__ static void CUDA_Func(BYTE* pDest, BYTE* pSrc1, BYTE* pSrc2, int nPos)
{
int i = nPos*512 + threadIdx.x * 1024; //各GPUスレッドが動き出す起点、縦にずれてるわけ
for (k = 0; k < 1024; k++) //各GPUスレッドは呼ばれると横向きに1024回ループする。内側ループをGPUスレッドで処理。
{
pDest[i + k] = ((int)pSrc1[i + k] + (int)pSrc2[i + k]) >> 1;
}
}
=====
ちなみにこれだといちいちグローバルメモリへのアクセスになるんで、
//Sharedにスクラッチコピー
//スクラッチコピー分だけループ処理
//SharedからGlobalに書き出し
するともっと速くなる
|
 |
