Un simulatore di onde con C# + Cudafy (C# -> CUDA o traduttore OpenCL) funziona alla grande, tranne che per l'esecuzione di OpenCL Versione CPU (driver Intel, 15 "MacBook Pro Retina i7 2.7GHz, GeForce 650M (Kepler, 384 core)) circa quattro volte più veloce della versione GPUCuda - CPU OpenCL 4x più veloce di OpenCL o CUDA versione GPU
(Questo accade se utilizzo la CL o la GPU CUDA . backend Le versioni OpenCL GPU e CUDA svolgono quasi identico)
per chiarire, per un problema di esempio:.
- OpenCL CPU 1200 Hz
- OpenCL GPU 320 Hz
- CUDA GPU - ~ 330 Hz
Sono in perdita per spiegare il motivo per cui la versione della CPU sarebbe più veloce rispetto alla GPU. In questo caso, il codice del kernel in esecuzione (nel caso CL) sulla CPU e GPU è identico. Seleziono la CPU o il dispositivo GPU durante l'inizializzazione, ma oltre a questo, tutto è identico.
Modifica
Ecco il codice C# che lancia uno dei noccioli. (Gli altri sono molto simili.)
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
);
}
Ed, ecco la funzione del kernel CUDA rilevanti generato da 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;
}
}
Solo per completezza, ecco la C# che viene utilizzato per generare il 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;
}
Ovviamente, il if in questo kernel è male, ma anche la pipeline di stallo risultante non dovrebbe causare un tale delta prestazioni estreme.
L'unica cosa che salta fuori di me è che sto usando uno schema di allocazione griglia/blocco lame - cioè, la griglia è la dimensione della matrice da aggiornare, e ciascun blocco è un thread. Sono sicuro che questo ha un impatto sulle prestazioni, ma non riesco a vederlo provocando un 1/4 della velocità del codice CL in esecuzione sulla CPU. ARGH!
Avete qualche esempio di codice che è possibile condividere? –
@EricBainville Sure - vuoi il C#, il kernel CUDA o CL, o cosa? (È un'applicazione semi-mid-size. Non voglio incollare 20k linee di codice in SO) –
Non vedo alcuna indicazione che il kernel di cuda stia utilizzando più di 1 thread per blocco (non c'è uso di 'threadIdx.x' o' threadIdx.y'). Inoltre il lancio sta specificando 1 thread per blocco. Ciò significa che circa il 97% delle capacità della GPU non è utilizzato. Non so molto su cudafy, quindi non so se hai il controllo su questo, ma non sono affatto sorpreso che il codice cuda non funzioni in modo impressionante veloce. –