有没有办法使用编译器在 AMD OpenCL 内核中展开循环？

Question

我正在尝试评估用于 AMD 和 Nvidia GPU 的 OpenCL 之间的性能差异。我有一个执行矩阵向量乘法的内核。我现在在两个不同的系统上运行内核，我的笔记本电脑有一个带有 Ubuntu 12.04 和 CUDA 4.0（包含 OpenCL 库和头文件）的 NVidia GT525m，另一个是带有 AMD Radeon HD7970 的台式机，再次带有 Ubuntu 12.04 和最新的 Catalyst 驱动程序。

在内核中，我有两个 #pragma unroll 语句，它们为 Nvidia OpenCL 实现（~6x）产生了很大的加速。然而，AMD OpenCL 版本不会产生任何加速。使用 AMD APP 内核分析器查看内核会出现未使用展开的错误，因为行程计数未知。 所以我的问题是，#pragma unroll 是否与 AMD OpenCL 一起工作，或者是否有替代方案（也许是我不知道的编译器标志）。我在下面包含了内核

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

相同的内核在两种实现中都产生了正确的结果，但是 #pragma unroll 命令对 AMD 没有任何作用（通过将它们注释掉来检查）。

Answer 1

它没有被记录，但它实际上应该与#pragma unroll 一起使用。您可以检查编译器日志以查看是否应用了展开？我不确定内核分析器是否使用与 OpenCL 运行时相同的编译器，您可能需要检查一下。

否则，如果您知道n 以 256 个块为单位，您可以手动展开，方法是在 256 个元素的块上循环一个循环，另一个循环在内部，固定大小为 256，这可能更容易展开。这样肯定能解决行程数静态不知道的问题。

但是，请记住，展开循环通常并不是什么大的胜利，因为您没有很多寄存器来缓存您的计算。循环展开增加的套准压力可能会导致套准溢出，这甚至更慢。您应该检查内核在 AMD 卡上的实际运行速度。较新的 NVIDIA OpenCL 编译器也可能不再从 unroll pragma 中受益。