Devo creare più OpenCL kernel per evitare istruzioni condizionali?

Question

In OpenCL, ho un kernel che ha bisogno di operare su dati complessi e reali. Ho potuto mettere un'istruzione condizionale in che chiama la linea giusta di codice per gestire questa situazione, o potrei avere due kernel che io chiamo e spingere l'istruzione condizionale fuori per il mio codice chiamante.

Questo è ovviamente un male per la manutenibilità, ma è significativo per le prestazioni?

Answer 1

Se è solo un'istruzione condizionale, nella mia esperienza la differenza di prestazioni è assolutamente trascurabili, almeno su hardware NVIDIA.

In sostanza, a patto che tutti (o quasi) di lavoro-articoli seguono lo stesso percorso di codice, che stai bene. Come il percorso di codice preso dipende da un argomento del kernel nel tuo caso, tutti i work-articoli seguono lo stesso percorso.

Answer 2

Dipende un po 'da dove il condizionale è. Codice per migliorare la leggibilità, poi prestazioni dopo aver misurato e trovò che si tratta di un problema

ad es. kernel_for_RGB_image e kernel_for_ABGR_image sembra un uso ragionevole, diversi kernel in modo efficace di srotolamento qualche profondo ciclo interno potrebbe essere un mal di testa di manutenzione più grande.

Answer 3

Penso che il modo migliore è quello di cercare in realtà di riferimento e due varianti. In alcuni casi, avere più blocchi condizionali compilati, anche se viene eseguito solo uno di loro, può portare a prestazioni peggiori. La ragione è GPRS (registri di uso generale):. Alloca compilatore altrettanti registri, e sono necessari per il caso peggiore

posso suggerire tale soluzione: hanno una sola funzione del kernel, ma con fase di compilazione condizionale:

__kernel void work()
{
#if VAR
    // one code
#else
    // another code
#endif
}

Allora avete bisogno di ricompilare il kernel con il set true / false a VAR quando si cambia condizione. Ovviamente, per il compilatore che non differisce da due kernel, ma per il mantenimento può essere meglio, se una parte del codice è lo stesso per questi kernel.