OpenCL Unrolling Loops优化

Posted 2020-12-19 zhuzhudong

　　可以控制离线编译器将kernel装换位硬件资源的方式。如果kernel中包含循环迭代，可以通过展开循环来提高性能。循环的展开减少了离线编译器执行的迭代次数，但代价是硬件资源消耗的增加。

　　如果有比较充分的硬件资源，直接在主循环中添加#progma unroll来展开循环。循环的展开会显著地改变离线编译器创建的计算单元的结构。

__kernel void example ( __global const int * restrict x,
                        __global int * restrict sum ) {
  int accum = 0;

  #pragma unroll
  for (size_t i = 0; i < 4; i++) {
    accum += x[i + get_global_id(0) * 4];
  }

  sum[get_global_id(0)] = accum;
}

　　#pragma unroll指令使离线编译器完全展开循环的四个迭代。 为了完成展开，离线编译器通过将加法运算数量增加三倍并加载四倍的数据来扩展pipeline。 移除循环后，计算单元将采用前馈结构。 结果，计算单元可以在初始加载操作和加法完成之后的每个时钟周期存储和的计算结果。 离线编译器通过合并四个加载操作来进一步优化此kernel，以便计算单元可以加载所有必需的输入数据从而在一个加载操作中计算出结果。

注意：

　　不要使用嵌套循环，尽可能地添加#progma unroll指令，实现大的单循环或展开内部循环。离线编译器不一定能轻易地展开循环，如果循环嵌套层数较多较复杂，会导致编译时间很长。

　　展开循环并合并全局内存中的加载操作，可使内核的硬件实现在每个时钟周期执行更多操作。 通常，用于提高OpenCL的kernel性能的方法应达到以下结果：

　　增加并行操作数

　　增加实现的内存带宽

　　增加内核可以在硬件中执行的每个时钟周期的操作数

　　而在以下情况下，离线编译器可能无法完全展开循环：

　　想要完全展开具有非常大数量迭代的存在数据依赖性的循环，kernel的硬件实现可能不适用于FPGA。

　　完全展开循环边界不是常数的循环。

　　循环由负责的控制flow组成，例如包含复杂数组索引或退出条件（在编译时未知）的循环。

　　要在这些情况下启用循环展开，指定#pragma unroll <N>指令，其中<N>是展开因子。 展开因子限制离线编译器展开的迭代次数。 例如，要阻止kernel中的循环展开，将指令#pragma unroll 1添加到该循环中。