解像度〜180X180の各入力画像と、解像度〜128 * 128の約10,000個のテンプレート画像を一致させるためのcudaプログラムを作成しています。目標は、リアルタイム パフォーマンスを達成することです。つまり、1 秒で 25 ～ 30 の入力画像 (それぞれに 10,000 個のテンプレートすべて) のテンプレート マッチングが行われます。
現在、私は次のアプローチを使用しています
- 実行時の I/O 操作を節約するために、すべてのテンプレートを GPU グローバル メモリにプリロードしました。
- 1 つのソース イメージとすべてのテンプレート イメージを一致させる単一のカーネルを作成し、すべての肯定的な一致の配列を返します。
- すべての操作を時間領域で実行します (FFT を使用しません)。その理由は、私は Radix-4 fft の実装を試みましたが、多くの中間グローバル読み取りと書き込みが必要であり、最終的に時間がかかります。
これまでのところ、1 つの入力画像から 10,000 のテンプレートに約 2 秒かかります。
私の質問は次のとおりです。
- このタスクがリアルタイムで達成可能かどうかを判断する方法はありますか? つまり、最大 FLOPS や I/O 帯域幅の制限などの助けを借りて
- GPU が最大限に活用されているかどうかを計算する方法は?
- パフォーマンスを改善する方法はありますか?
マシンスペック：[i7-4770、8GB、GTX-680]
現在のカーネル コードの説明:
- すべてのテンプレート イメージ [サイズは RGB で約 128X128] は、GPU メモリに個別に読み込まれます。アイデアは、実行時の操作中に I/O を節約することです。
- すべての入力画像はテクスチャ メモリにロードされます。その理由は、テクスチャが 2D アドレッシングに適したオプションであるためです。
- すべての「ブロック」には 1024 のスレッドがあります。
- 各スレッドは各出力ピクセルの値を計算します。出力のサイズは [31X31 = 961 ピクセル] です。
- 起動されるブロックの数は、一致するテンプレート イメージの数と同じです。
カーネル コード:
__global__ void cudaMatchTemplate(TemplateArray *templates, uchar *Match)
{
int global = blockIdx.x*blockDim.x + threadIdx.x;
__shared__ int idx[TEMPLATE_MATCH_DIM];
__shared__ float out_shared[TEMPLATE_MATCH_DIM];
//halving the template size....
int rows = (templates[blockIdx.x].nHeight)/2;
int cols = (templates[blockIdx.x].nWidth)/2;
int fullCol = templates[blockIdx.x].nWidth;
int x = templates[blockIdx.x].nMatchLeft;
int y = templates[blockIdx.x].nMatchTop;
int offset_y = (threadIdx.x/TEMPLATE_MATCH_SIZE);
int offset_x = (threadIdx.x - offset_y*TEMPLATE_MATCH_SIZE);
// *************** Performing match in time domain *****************************//
int sum = 0;
float temp;
int idxXFactor = 3*(2*(offset_x) + x);
int idxYFactor = 2*(offset_y) + y ;
for (int i = 0; i < rows; i++)
{
int I=3*i*fullCol;
int sourceIdxY = idxYFactor + 2*i;
for (int j = 0; j < cols; j++)
{
int J=3*j;
int sourceIdxX = idxXFactor + 2*J;
int templateIdx = 2*I+2*J;
//**** R *****//
temp = float(tex2D(SourceImgColorTex,sourceIdxX,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx]);
sum = sum + temp*temp;
//**** G *****//
temp = float(tex2D(SourceImgColorTex,sourceIdxX+1,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +1]);
sum = sum + temp*temp;
//**** B *****//
temp = float(tex2D(SourceImgColorTex,sourceIdxX+2,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +2]);
sum = sum + temp*temp;
}
}
__syncthreads();
//placing all values in shared memory for comparison.
if(threadIdx.x < TEMPLATE_MATCH_DIM)
{
idx[threadIdx.x] = threadIdx.x;
out_shared[threadIdx.x] = sum;
}
__syncthreads();
// //computing the Min location.....//
#pragma unroll
for(int s=512; s>0; s>>=1)
{
if ((threadIdx.x < s) &&((threadIdx.x + s)<TEMPLATE_MATCH_DIM))
{
idx[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? idx[threadIdx.x] : idx[threadIdx.x + s];
out_shared[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? out_shared[threadIdx.x] : out_shared[threadIdx.x + s];
}
}
__syncthreads();
if(threadIdx.x <1)
{
int half_Margin = MARGIN_FOR_TEMPLATE_MATCH/2;
int matchY = idx[0]/TEMPLATE_MATCH_SIZE ;
int matchX = idx[0] - matchY * TEMPLATE_MATCH_SIZE;
int diff = absolute(half_Margin - matchX) + absolute(half_Margin - matchY);
if(diff < THRESHOLD)
{
Match[blockIdx.x] = 1;
}
else
Match[blockIdx.x] = 0;
}
}