Gibt es eine Möglichkeit, Schleifen in einem AMD OpenCL-Kernel mit dem Compiler zu lösen?

Question

Ich versuche, die Leistungsunterschiede zwischen OpenCL für AMD und Nvidia GPUs zu bewerten. Ich habe einen Kernel, der Matrix-Vektor-Multiplikation durchführt. Ich betreibe den Kernel zu diesem Zeitpunkt auf zwei verschiedenen Systemen, meinem Laptop mit NVidia GT525m mit Ubuntu 12.04 und CUDA 4.0 (der die OpenCL-Bibliotheken und Header enthält) und der andere ist ein Desktop mit einer AMD Radeon HD7970 wieder mit Ubuntu 12.04 und die neuesten Catalyst-Treiber.

Im Kernel habe ich zwei #pragma unroll Anweisungen, die eine große Beschleunigung für die Nvidia OpenCL-Implementierung (~ 6x) erzeugen. Allerdings bringt die AMD OpenCL-Version keine Beschleunigung. Betrachtet man den Kernel mit dem AMD APP-Kernel-Analyzer, tritt der Fehler auf, dass die Unroll-Funktion nicht verwendet wird, da die Anzahl der Trips nicht bekannt ist. Also meine Frage ist, funktioniert #pragma unroll mit AMD OpenCL oder gibt es eine Alternative (vielleicht ein Compiler-Flag, das ich nicht weiß). Ich habe den Kernel enthalten unter

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

Diese gleichen Kernel korrekte Ergebnisse in beiden Implementierungen erzeugt aber die Pragma entrollen Befehle tun nichts für die AMD (geprüft von ihnen heraus zu kommentieren).

Answer 1

Es ist nicht dokumentiert, aber es sollte tatsächlich mit #pragma unroll arbeiten. Können Sie das Compiler-Protokoll überprüfen, um zu sehen, ob das Ausrollen angewendet wird? Ich bin mir nicht sicher, ob der Kernel-Analyzer den gleichen Compiler wie die OpenCL-Laufzeit verwendet, was Sie vielleicht überprüfen möchten.

Sonst, wenn Sie wissen, dass n in 256 Stücken kommt, können Sie manuell abrollen, indem Sie eine Schleife über Blöcke von 256 Elementen und eine andere innerhalb einer festen Größe von 256 haben, die einfacher zu entrollen sein könnte. Dies wird sicherlich das Problem lösen, dass die Reisezählung nicht statisch bekannt ist.

Beachten Sie jedoch, dass das Abwickeln einer Schleife normalerweise nicht so groß ist, da Sie nicht viele Register haben, um Ihre Berechnungen zwischenzuspeichern. Der erhöhte Regeldruck beim Schleifenabrollen könnte zu einem Überlaufen des Registers führen, das noch langsamer ist. Sie sollten überprüfen, wie schnell der Kernel tatsächlich auf der AMD-Karte ist. Ein neuerer NVIDIA OpenCL-Compiler kann möglicherweise auch nicht mehr vom Pragma abrollen.