次はcudaでプログラミングしています:
- 画像値 (d_Data) を含むマトリックスを取得します
- 行列を共有メモリ (タイル) にコピーします
- 基本的なピクセル差 (pixel(i)-pixel(i+1)) (d_diff) を取得します
- 差分が特定の値、たとえば 0 の場合、行列 (d_A) の d_diff で数値 0 が見つかった各位置に数値 1 が設定されます。これは、元の差分行列で 0 の頻度を取得できるようにするためです。
- 並列累積合計。
- 周波数結果は周波数ベクトルに移動します。
順を追って確認すると、すべてが累積合計まで期待どおりに進んでいます。コードを起動すると、ソフトウェアで計算された値は 104347 ですが、CUDA から nan の結果が得られることもあれば、2425 などの任意の数値が得られることもあります。非常に奇妙なのは、カーネルを 20 回または 30 回実行し続けると、値が次のようになることです。予想される 104347 :S.
私は各マトリックスに使用しています:
h_Data = (float *)malloc(data_size);
h_diff = (float *)malloc(data_size);
h_A = (float *)malloc(data_size);
と
cudaFree(d_A);
cudaFree(d_diff);
cudaFree(d_Av);
そのため、十分な回数実行するとコードが正しい結果に近づく理由がわかりません。ちなみに、正しい値になると、何度コードを実行しても動かなくなります。
コード:
__global__ void spam(float *d_Data, float *d_diff, float *d_A, int dw, int dh, float *d_Av){
long bx = blockIdx.x; long by = blockIdx.y;
long tx = threadIdx.x; long ty = threadIdx.y;
// Identify the row and column of the Pd element to work on
long Row = by * TILE_WIDTH + ty;
long Col = bx * TILE_WIDTH + tx;
long tid = Row*dw+Col;
long i=512*512;
long r = MASK_DIM/2;
long s = 0;
__shared__ int tile[BLOCK_WIDTH][BLOCK_WIDTH];
for (int k=0; k<=8; k++)
d_Av[k]=0;
if(tid < dw*dh)
{
// to shared memory.
tile[ty + r][tx + r]=d_Data[Row*dw+Col];
if (Col-r >=0) tile[ty + r] [tx] = d_Data[Row*dw+Col-r];
if (Col+r <dw) tile[ty + r] [tx + 2*r] = d_Data[Row*dw+Col+r];
if (Row-r >=0) tile[ty] [tx + r] = d_Data[(Row - r)*dw + Col];
if (Row+r <dw) tile[ty + 2*r][tx + r] = d_Data[(Row + r)*dw + Col];
if (Row - r >= 0 && Col - r >= 0) tile[ty] [tx] = d_Data[(Row-r)*dw+Col-r];
if(Row - r >= 0 && Col + r < dw) tile[ty] [tx + 2*r] = d_Data[(Row-r)*dw+Col+r];
if (Row + r < dw && Col - r >= 0) tile[ty + 2*r][tx] = d_Data[(Row+r)*dw+Col-r];
if(Row + r <dw && Col + r < dw) tile[ty + 2*r][tx + 2*r] = d_Data[(Row-r)*dw+Col+r];
//Calculates the difference matrix
d_diff[tid] = (tile[ty + r][tx +r] - tile[ty + r][tx + r + 1]);
d_A[tid]=0;
//Set a 1 in each position in d_A where 0 was found in d_diff.
if (d_diff[tid] == 0)
{ d_A[tid]=1;}
__syncthreads();
//cumulative sum to get the frecuency of value 0 in d_diff. // The error is HERE
for (s = (i/2); s>=1; s=s/2) {
if (tid < s)
{ d_A[tid] += d_A[tid+s];
}
}
// set the frequency value in frequencies vector.
d_Av[0] = d_A[0];
}} // END IF tid < dw*dh
どんなアイデアでも大歓迎です:D