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

Question

我正在尝试评估 OpenCL for AMD 和 Nvidia GPU 之间的性能差异。我有一个执行矩阵向量乘法的内核。我目前正在两个不同的系统上运行内核，我的笔记本电脑有一个带有 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 编译器也可能不再从展开编译指示中受益。