Cuda - OpenCL CPU 4x szybciej niż wersja GPU OpenCL lub CUDA

Question

Symulator fal, nad którym pracowałem z C# + Cudafy (C# -> CUDA lub OpenCL translator) działa świetnie, z wyjątkiem faktu, że uruchamianie OpenCL Wersja CPU (sterownik Intel, 15-calowy MacBook Pro Retina i7 2.7GHz, GeForce 650M (Kepler, 384 rdzenie)) jest mniej więcej cztery razy szybsza od wersji GPU. . backend wersje OpenCL graficznego i CUDA wykonać prawie identycznie)

celu wyjaśnienia, do problemu próbki.

• OpenCL CPU 1200 Hz
• OpenCL GPU 320 Hz
• CUDA GPU - ~ 330 Hz

Jestem w rozterce wyjaśnić, dlaczego wersja CPU byłoby szybciej niż GPU. W tym przypadku kod jądra, który wykonuje (w przypadku CL) na CPU i GPU jest identyczny. Wybieram CPU lub GPU podczas inicjalizacji, ale poza tym wszystko jest identyczne.

Edit

Oto kod C#, który uruchamia jedną z jądrami. (Pozostałe są bardzo podobne.)

public override void UpdateEz(Source source, float Time, float ca, float cb) 
    { 
     var blockSize = new dim3(1); 
     var gridSize = new dim3(_gpuEz.Field.GetLength(0),_gpuEz.Field.GetLength(1)); 

     Gpu.Launch(gridSize, blockSize) 
      .CudaUpdateEz(
       Time 
       , ca 
       , cb 
       , source.Position.X 
       , source.Position.Y 
       , source.Value 
       , _gpuHx.Field 
       , _gpuHy.Field 
       , _gpuEz.Field 
      ); 

    } 

A oto właściwa funkcja jądra CUDA generowane przez Cudafy:

extern "C" __global__ void CudaUpdateEz(float time, float ca, float cb, int sourceX, int sourceY, float sourceValue, float* hx, int hxLen0, int hxLen1, float* hy, int hyLen0, int hyLen1, float* ez, int ezLen0, int ezLen1) 
{ 
    int x = blockIdx.x; 
    int y = blockIdx.y; 
    if (x > 0 && x < ezLen0 - 1 && y > 0 && y < ezLen1 - 1) 
    { 
     ez[(x) * ezLen1 + (y)] = ca * ez[(x) * ezLen1 + (y)] + cb * (hy[(x) * hyLen1 + (y)] - hy[(x - 1) * hyLen1 + (y)]) - cb * (hx[(x) * hxLen1 + (y)] - hx[(x) * hxLen1 + (y - 1)]); 
    } 
    if (x == sourceX && y == sourceY) 
    { 
     ez[(x) * ezLen1 + (y)] += sourceValue; 
    } 
} 

Tylko dla kompletności, tutaj jest C#, który jest używany do generowania CUDA:

[Cudafy] 
    public static void CudaUpdateEz(
     GThread thread 
     , float time 
     , float ca 
     , float cb 
     , int sourceX 
     , int sourceY 
     , float sourceValue 
     , float[,] hx 
     , float[,] hy 
     , float[,] ez 
     ) 
    { 
     var i = thread.blockIdx.x; 
     var j = thread.blockIdx.y; 

     if (i > 0 && i < ez.GetLength(0) - 1 && j > 0 && j < ez.GetLength(1) - 1) 
      ez[i, j] = 
       ca * ez[i, j] 
       + 
       cb * (hy[i, j] - hy[i - 1, j]) 
       - 
       cb * (hx[i, j] - hx[i, j - 1]) 
       ; 

     if (i == sourceX && j == sourceY) 
      ez[i, j] += sourceValue; 
    } 

Oczywiście, if w tym jądrze jest złe, ale nawet wynikowe przeciągnięcie rurociągu nie powinno powodować tak ekstremalnej wydajności.

Jedyną inną rzeczą, która wyskakuje na mnie jest to, że używam lame grid/block scheme scheme - tj. Siatka to rozmiar tablicy do aktualizacji, a każdy blok to jeden wątek. Jestem pewien, że ma to pewien wpływ na wydajność, ale nie widzę powodując, że jest to 1/4 szybkości kodu CL działającego na procesorze. ARGH!

Answer 1

Odpowiedzi na to pytanie, aby uzyskać go z listy bez odpowiedzi.

Opublikowany kod oznacza, że ​​uruchomienie jądra określa wątek z wątkiem 1 (aktywnym). Nie jest to sposób na zapisanie szybkiego kodu GPU, ponieważ pozostawia on większość możliwości GPU bezczynności.

Typowe rozmiary bloków wątku powinny wynosić co najmniej 128 wątków na blok, a wyższe często są lepsze, w wielokrotnościach 32, do limitu 512 lub 1024 na blok, w zależności od GPU.

Procesor graficzny "lubi", aby ukryć opóźnienie, wykonując wiele zadań równoległych "dostępnych". Podanie więcej wątków na blok pomaga temu celowi. (Możliwa jest również odpowiednia liczba bloków wątku w siatce.)

Ponadto GPU wykonuje wątki w grupach po 32.Podanie tylko jednego wątku na blok lub nie wielokrotnego z 32 spowoduje pozostawienie wolnych gniazd wykonawczych w każdym bloku wątków, który zostanie wykonany. 1 wątek na blok jest szczególnie zły.