Я реализовал простое ядро, которое является своего рода сверткой. Я измерял это на NVIDIA GT 240. Это заняло 70 мс при написании на CUDA и 100 мс при написании на OpenCL. Хорошо, я думал, компилятор 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 мс.
Любые замечания по поводу моего кода также приветствуются.