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

Question

我正在尝试评估OpenCL for AMD和Nvidia GPU之间的性能差异.我有一个执行矩阵向量乘法的内核.我现在在两个不同的系统上运行内核,我的笔记本电脑有一个带Ubuntu 12.04的NVidia GT525m和CUDA 4.0(包含OpenCL库和标题),另一个是带有Ubuntu的AMD Radeon HD7970的桌面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运行时相同的编译器,您可能需要检查.

否则,如果你知道n256块的块,你可以手动展开一个循环超过256个元素的块,另一个循环在256的固定大小,这可能更容易展开.这肯定会解决旅行计数静态不知道的问题.

但是,请记住,展开循环通常不是一个很大的胜利,因为你没有很多寄存器来缓存你的计算.来自循环展开的寄存器压力增加可能导致寄存器溢出,这甚至更慢.您应该检查内核实际上在AMD卡上的速度.较新的NVIDIA OpenCL编译器也可能不再受益于unroll编译指示.