Я реализовал простое ядро, которое является своего рода сверткой. Я измерил его на NVIDIA GT 240. При записи на CUDA потребовалось 70 мс, а при записи на OpenCL — 100 мс. Ладно, подумал я, компилятор NVIDIA лучше оптимизирован под CUDA (или я что-то не так делаю). Мне нужно запустить его на графических процессорах AMD, поэтому я перешел на AMD APP SDK. Точно такой же код ядра.
Я провел два теста, и их результаты меня обескуражили: 200 мс на HD 6670 и 70 мс на HD 5850 (столько же, сколько у GT 240 + CUDA). И меня очень интересуют причины такого странного поведения.
Все проекты были построены на VS2010 с использованием настроек примеров проектов NVIDIA и AMD соответственно.
Пожалуйста, не считайте мой пост рекламой NVIDIA. Я прекрасно понимаю, что HD 5850 мощнее GT 240. Единственное, что хотелось бы знать, так это почему такая разница и как решить проблему.
Обновлять. Ниже приведен код ядра, который ищет 6 шаблонных изображений одинакового размера в базовом. Каждый пиксель базового изображения рассматривается как возможный источник одного из шаблонов и обрабатывается отдельным потоком. Ядро сравнивает значения R, G, B каждого пикселя базового изображения и эталонного, и если хотя бы одно отличие превышает параметр diff, соответствующий пиксель считается несовпадающим. Если количество несовпадающих пикселей меньше maxNonmatchQt, соответствующий шаблон срабатывает.
__constant int tOffset = 8196; // one template size in memory (in bytes)
__kernel void matchImage6( __global unsigned char* image, // pointer to the base image
int imgWidth, // base image width
int imgHeight, // base image height
int imgPitch, // base image pitch (in bytes)
int imgBpp, // base image bytes (!) per pixel
__constant unsigned char* templates, // pointer to the array of templates
int tWidth, // templates width (the same for all)
int tHeight, // templates height (the same for all)
int tPitch, // templates pitch (in bytes, the same for all)
int tBpp, // templates bytes (!) per pixel (the same for all)
int diff, // max allowed difference of intensity
int maxNonmatchQt, // max number of nonmatched pixels
__global int* result, // results
) {
int x0 = (int)get_global_id(0);
int y0 = (int)get_global_id(1);
if( x0 + tWidth > imgWidth || y0 + tHeight > imgHeight)
return;
int nonmatchQt[] = {0, 0, 0, 0, 0, 0};
for( int y = 0; y < tHeight; y++) {
int ind = y * tPitch;
int baseImgInd = (y0 + y) * imgPitch + x0 * imgBpp;
for( int x = 0; x < tWidth; x++) {
unsigned char c0 = image[baseImgInd];
unsigned char c1 = image[baseImgInd + 1];
unsigned char c2 = image[baseImgInd + 2];
for( int i = 0; i < 6; i++)
if( abs( c0 - templates[i * tOffset + ind]) > diff ||
abs( c1 - templates[i * tOffset + ind + 1]) > diff ||
abs( c2 - templates[i * tOffset + ind + 2]) > diff)
nonmatchQt[i]++;
ind += tBpp;
baseImgInd += imgBpp;
}
if( nonmatchQt[0] > maxNonmatchQt && nonmatchQt[1] > maxNonmatchQt && nonmatchQt[2] > maxNonmatchQt && nonmatchQt[3] > maxNonmatchQt && nonmatchQt[4] > maxNonmatchQt && nonmatchQt[5] > maxNonmatchQt)
return;
}
for( int i = 0; i < 6; i++)
if( nonmatchQt[i] < maxNonmatchQt) {
unsigned int pos = atom_inc( &result[0]) * 3;
result[pos + 1] = i;
result[pos + 2] = x0;
result[pos + 3] = y0;
}
}
Конфигурация запуска ядра: Глобальный рабочий размер = (1900, 1200) Локальный рабочий размер = (32, 8) для AMD и (32, 16) для NVIDIA.
Время выполнения: HD 5850 — 69 мс, HD 6670 — 200 мс, GT 240 — 100 мс.
Любые замечания о моем коде также высоко ценятся.
