Czy istnieje sposób na rozwinięcie pętli w jądrze AMD OpenCL z kompilatorem?

Question

Próbuję ocenić różnice wydajności między OpenCL dla procesorów graficznych AMD i Nvidia. Mam jądro, które wykonuje mnożenie macierzy-wektora. Obecnie uruchamiam jądro w dwóch różnych systemach, mój laptop z NVIDIA GT525m z Ubuntu 12.04 i CUDA 4.0 (który zawiera biblioteki i nagłówki OpenCL), a drugi to komputer z AMD Radeon HD7970 ponownie z Ubuntu 12.04 i najnowsze sterowniki Catalyst.

W jądrze mam dwie instrukcje #pragma unroll, które dają duże przyspieszenie dla implementacji Nvidia OpenCL (~ 6x). Jednak wersja AMD OpenCL nie generuje żadnego przyspieszenia. Patrzenie na jądro za pomocą analizatora jądra APP AMD daje błąd, że rozwijanie nie jest używane, ponieważ licznik podróży nie jest znany. Moje pytanie brzmi, czy #pragma unroll współpracuje z AMD OpenCL, czy istnieje alternatywa (być może flaga kompilatora, której nie znam). Podaję jądra poniżej

__kernel void mvKernel(__global float* a, const __global float* x, __global float* y, int m, int n) 
{ 
    float sum = 0.0f; 
    __global float* A; 
    int i; 
    int j = 0; 
    int indx = get_global_id(0); 
    __local float xs[12000]; 
#pragma unroll 
    for(i = get_local_id(0); i < n; i+= get_local_size(0)) { 
     xs[i] = x[i]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 
    A = &a[indx]; 
#pragma unroll 256 
    for(i = 0; i < n; i++) { 
     sum += xs[i] * A[j]; 
     j += m; 
    } 
    y[indx] = sum; 
} 

To samo jądro produkuje poprawne wyniki w obu implementacjach ale polecenia #pragma UNROLL nic nie robią dla AMD (sprawdzony przez komentując je).

Answer 1

To nie jest dokumentowane, ale powinno działać z #pragma unroll. Czy możesz sprawdzić dziennik kompilatora, aby sprawdzić, czy rozwijanie jest stosowane? Nie jestem pewien, czy analizator jądra używa tego samego kompilatora, co środowisko uruchomieniowe OpenCL, możesz to sprawdzić.

W przeciwnym razie, jeśli wiesz, że n ma porcje 256, możesz rozwinąć ręcznie, mając jedną pętlę ponad blokami składającymi się z 256 elementów, a drugą ze stałym rozmiarem 256, co może być łatwiejsze do rozwinięcia. To na pewno rozwiąże problem polegający na tym, że liczba wyzwoleń nie jest znana statycznie.

Pamiętaj jednak, że rozwijanie pętli zwykle nie jest tak dużym zwycięstwem, ponieważ nie masz wielu rejestrów do buforowania obliczeń. Zwiększone ciśnienie rejestrowe z rozwijania pętli może prowadzić do zarejestrowania rozlewania, które jest nawet wolniejsze. Powinieneś sprawdzić, jak szybko jądro faktycznie znajduje się na karcie AMD. Nowszy kompilator NVIDIA OpenCL może również nie zyskać więcej na praglecie rozwijania.