OpenCL AMD vs NVIDIA performance

Question

Zaimplementowałem proste jądro, które jest rodzajem splotu. Zmierzyłem to na NVIDIA GT 240. Zajęło to 70 ms, gdy napisałem na CUDA i 100 ms, gdy napisałem na OpenCL. Ok, pomyślałem, kompilator NVIDIA jest lepiej zoptymalizowany dla CUDA (lub robię coś źle). Muszę go uruchomić na procesorach graficznych AMD, więc migrowałem do SDK aplikacji AMD. Dokładnie ten sam kod jądra.

Zrobiłem dwa testy i ich wyniki były dla mnie zniechęcające: 200 ms w HD 6670 i 70 ms w HD 5850 (w tym samym czasie co w GT 240 + CUDA). Jestem bardzo zainteresowany przyczynami tak dziwnego zachowania.

Wszystkie projekty zostały zbudowane na VS2010 przy użyciu ustawień z przykładowych projektów odpowiednio NVIDIA i AMD.

Proszę, nie traktuj mojego postu jako reklamy NVIDIA. Rozumiem, że HD 5850 jest silniejszy niż GT 240. Jedyne, co chcę wiedzieć, to dlaczego taka jest różnica i jak rozwiązać problem.

Aktualizacja. Poniżej znajduje się kod jądra, który wyszukuje 6 obrazów o jednakowej wielkości w szablonie. Każdy piksel obrazu podstawowego jest uważany za możliwe pochodzenie jednego z szablonów i jest przetwarzany przez osobny wątek. Jądro porównuje wartości R, G, B każdego piksela obrazu podstawowego i szablonu pierwszego, a jeśli co najmniej jedna różnica przekracza parametr diff, odpowiadający piksel jest zliczany niedopasowany. Jeśli liczba niedopasowanych pikseli jest mniejsza niż maxNonmatchQt, odpowiedni szablon zostanie trafiony.

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

Kernel konfiguracja run: Globalna wielkość praca = (1900, 1200) lokalny rozmiar praca = (32, 8) i AMD (32, 16) dla NVIDIA.

Czas realizacji: HD 5850 - 69 ms, HD 6670 - 200 ms, GT 240 - 100 ms.

Wszelkie uwagi na temat mojego kodu są również bardzo cenne.

Answer 1

Nie może być dokładnej idealnej odpowiedzi na to. Wydajność OpenCL zależy od wielu parametrów. Liczba dostępu do pamięci globalnej, efektywność kodu itp. Ponadto bardzo trudne porównanie dwóch urządzeń, ponieważ mogą one mieć różne lokalne, globalne, stałe wspomnienia. Liczba rdzeni, częstotliwość, przepustowość pamięci, a co ważniejsze architektura sprzętu itp.

Każdy sprzęt zapewnia własne zwiększenie wydajności, na przykład native_ firmy NVIDIA. Musisz więc dowiedzieć się więcej o sprzęcie, na którym pracujesz, co może faktycznie zadziałać. Ale osobiście polecam, aby nie używać takich specyficznych optymalizacji sprzętu, które mogłyby wpłynąć na elastyczność kodu.

Można również znaleźć niektóre opublikowane artykuły, które pokazują, że wydajność CUDA jest znacznie lepsza niż wydajność OpenCL na tym samym sprzęcie NVIDIA.

Zawsze lepiej napisać kod, który zapewnia dużą elastyczność, a nie optymalizację urządzenia.

Answer 2

Różnica czasów wykonania jest spowodowana przez kompilatory. Twój kod może być łatwo wektoryzowany. Rozważ obraz i szablony jako tablice typu wektorowego char4 (czwarta współrzędna każdego wektora char4 ma zawsze wartość 0).Zamiast 3 pamięć brzmi:

unsigned char c0 = image[baseImgInd]; 
unsigned char c1 = image[baseImgInd + 1]; 
unsigned char c2 = image[baseImgInd + 2]; 

użycie tylko jednego:

unsigned char4 c = image[baseImgInd]; 

zamiast wielkogabarytowych jeżeli:

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]++; 

stosowanie szybko:

unsigned char4 t = templates[i * tOffset + ind]; 
    nonmatchQt[i] += any(abs_diff(c,t)>diff); 

W ten sposób zwiększyć wydajność kodu do 3 razy (jeśli komp ler nie wektoryzuje kodu samoczynnie). Przypuszczam, że kompilator AMD OpenCL nie pozwala na wektoryzację i inne optymalizacje. Z mojego doświadczenia OpenCL na GPU NVIDIA zazwyczaj można zrobić szybciej niż CUDA, ponieważ jest on bardziej niski.