Gibt es eine Möglichkeit zu entrollen von Schleifen in einem AMD-OpenCL-kernel mit dem compiler?

Ich versuche zu bewerten, die performance-Unterschiede zwischen OpenCL für AMD-und Nvidia-Grafikkarten. Ich habe einen kernel, der führt eine matrix-Vektor-Multiplikation. Bei mir läuft der kernel auf zwei verschiedenen Systemen in den Momenten, mein laptop hat eine NVidia GT525m mit Ubuntu 12.04 und CUDA 4.0 (enthält die OpenCL-Bibliotheken und Header-Dateien) und der andere ist ein desktop mit einem AMD Radeon HD7970 wieder mit Ubuntu 12.04 und dem neuesten Catalyst-Treiber.

Im kernel habe ich zwei #pragma unroll Aussagen, die produzieren eine große speed-up für den Nvidia-OpenCL-Implementierung (~6x). Aber die AMD OpenCL-version produziert kein speedup. Blick auf den kernel mit dem AMD-APP-kernel-analyzer gibt den Fehler, dass die Rollen nicht verwendet, weil die Reise zu zählen, ist nicht bekannt.
Also meine Frage ist, hat #pragma unroll Arbeit mit AMD OpenCL oder gibt es eine alternative (vielleicht ein compiler-flag, dass ich bewusst bin). Ich habe den kernel 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 erzeugt die richtigen Ergebnisse in beiden Implementierungen, aber die #pragma-Rollen-Befehle tun nichts für die AMD (geprüft, kommentieren Sie Sie aus).