Как определить производительность CUDA GPU?

Я пишу программу cuda для сопоставления каждого входного изображения с разрешением ~ 180X180, с около 10000 шаблонных изображений с разрешением ~ 128 * 128. Цель состоит в достижении производительности в реальном времени, т. Е. Сопоставление шаблонов 25 ~ 30 входных изображений (каждое со всеми 10000 шаблонов) за 1 секунду.

В настоящее время я использую следующий подход

1. Предварительно загружены все шаблоны в глобальную память графического процессора для сохранения операций ввода-вывода во время выполнения.
2. Создано одно ядро ​​для сопоставления одного исходного изображения со всеми шаблонными изображениями и возврата массива для всех положительных совпадений.
3. Выполнение всех операций во временной области (без использования БПФ). причина в том, что я попробовал реализацию Radix-4 fft, но она требует много промежуточных глобальных операций чтения и записи, что в итоге занимает больше времени.

Пока для 1 входного изображения в 10000 шаблонов требуется около 2 секунд.

Мои вопросы:

1. Есть ли способ определить, достижима ли эта задача в реальном времени или нет? Я имею в виду максимальные ограничения пропускной способности FLOPS и I / O, т. Е.
2. Как рассчитать, если GPU полностью используется на максимуме?
3. Возможные способы улучшения производительности?

Технические характеристики машины: [i7-4770, 8 ГБ, GTX-680]

Объяснение текущего кода ядра:

1. все изображения шаблонов (размер около 128X128 в RGB) загружаются в память графического процессора. Идея состоит в том, чтобы сохранить ввод / вывод во время работы.
2. Каждое входное изображение загружается в текстурную память, поскольку текстура является хорошим вариантом для 2D-адресации.
3. Каждый «Блок» имеет 1024 темы.
4. Каждый поток вычисляет значение для каждого выходного пикселя, размер вывода составляет [31X31 = 961 пикселей].
5. Количество запущенных блоков равно количеству сопоставляемых шаблонных изображений.

Код ядра:

__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;

}
}