C'è un modo per srotolare i loop in un kernel OpenCL di AMD con il compilatore?

Question

Sto provando a valutare le differenze di prestazioni tra OpenCL per le GPU AMD e Nvidia. Ho un kernel che esegue la moltiplicazione di matrice-vettore. Sto eseguendo il kernel su due sistemi diversi in questo momento, il mio laptop che ha una NVidia GT525m con Ubuntu 12.04 e CUDA 4.0 (che contiene le librerie e le intestazioni OpenCL) e l'altro è un desktop con una AMD Radeon HD7970 di nuovo con Ubuntu 12.04 e gli ultimi driver Catalyst.

Nel kernel ho due istruzioni #pragma unroll che producono un grande aumento di velocità per l'implementazione Nvidia OpenCL (~ 6x). Tuttavia, la versione OpenCL di AMD non produce alcuna accelerazione. Guardando il kernel con l'analizzatore del kernel APP AMD si dà l'errore che il srotolamento non viene utilizzato perché il conteggio del viaggio non è noto. Quindi la mia domanda è: #pragma unroll funziona con AMD OpenCL o c'è un'alternativa (forse un flag del compilatore di cui non sono a conoscenza). Ho incluso il kernel sotto

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

Questo stesso kernel produce risultati corretti in entrambe le implementazioni ma i comandi #pragma srotolano non fare nulla per l'AMD (controllato da commentandoli).

Answer 1

Non è documentato, ma in realtà dovrebbe funzionare con #pragma unroll. Puoi controllare il log del compilatore per vedere se il srotolamento è applicato? Non sono sicuro che l'analizzatore del kernel utilizzi lo stesso compilatore del runtime OpenCL, potresti voler controllare.

Altrimenti, se si sa che n arriva in blocchi di 256, è possibile eseguire manualmente lo sbroglio avendo un ciclo su blocchi di 256 elementi e un altro all'interno con una dimensione fissa di 256, che potrebbe essere più facile da srotolare. Questo risolverà sicuramente il problema che il conteggio del viaggio non è noto staticamente.

Tuttavia, tenere a mente lo srotolamento di un ciclo di solito non è una grande vittoria, in quanto non si dispone di molti registri per memorizzare il calcolo. La maggiore pressione del registro dovuta allo srotolamento del circuito potrebbe causare la fuoriuscita del registro, che è ancora più lento. Dovresti controllare quanto velocemente il kernel si trova effettivamente sulla scheda AMD. Un nuovo compilatore NVIDIA OpenCL potrebbe anche non beneficiare più del pragma di srotolamento.