OpenCL AMD против производительности NVIDIA

Я реализовал простое ядро, которое представляет собой своего рода свертку. Я измерял его на 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 мс.

Любые замечания по поводу моего кода также приветствуются.