4.2. 展开循环

Intel® FPGA SDK for OpenCL™ Pro Edition: 最佳实践实践指南

下载 PDF

ID 683521

日期 9/26/2022

版本 22.3

Public

仅对英特尔可见 — GUID： mwh1391807501939

Ixiasoft

查看详细信息

4.2. 展开循环

您可以控制 Intel® FPGA SDK for OpenCL™ Offline Compiler将OpenCL™内核描述转换成硬件资源的方式。如果您的OpenCL内核包含循环迭代，请展开循环来提高性能。循环展开减少了离线编译器执行的迭代次数，但代价是增加了硬件资源消耗。

提示：对于 Intel® oneAPI DPC++/C++ Compiler的具体细节，请参阅 Intel® oneAPI Toolkits的FPGA优化指南中的展开循环部分。

请考虑并行应用程序的OpenCL代码，其中每个work-item负责阵列中4个单元的累加：

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

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

   sum[get_global_id(0)] = accum;
}

请注意内核中发生的以下3个主要操作：

从输入X进行加载操作
累加（Accumulation）
将操作存储到输出sum

离线编译器根据OpenCL内核代码的数据流语义（semantic）将这些操作以流水线形式排列。例如，离线编译器通过将结果从流水线末端转发到流水线顶部来实现循环，具体取决于循环退出条件。

OpenCL内核在每个时钟周期执行每个work-item的一次循环迭代。有了足够的硬件资源，您可以展开循环来提高内核性能，从而减少内核执行的迭代次数。要展开循环，请将#pragma 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指令使离线编译器完全展开循环的四个迭代。为了完成展开，离线编译器将加法操作的数量增加三倍并加载四倍的数据来扩展流水线。随着循环的移除，计算单元采取前馈（feed-forward）结构。从而，计算单元可以在初始加载操作和加法完成后的每个时钟周期存储sum单元。离线编译器将4个加载操进行合并以进一步优化该内核，以便计算单元能够加载所有必要的输入数据加载并通过一次加载操作来计算结果。

警告：

避免嵌套式循环结构。反而应尽可能通过添加#pragma unroll指令来实现大型但循环或展开内部循环。

例如，如果您编译的内核具有高度嵌套的循环结构，其中每个循环都包含一个#pragma unroll指令，您可能会经历漫长的编译时间。 Intel® FPGA SDK for OpenCL™ Offline Compiler可能无法满足调度，因为它不能轻松展开该嵌套循环结构，从而导致较高的II。该情况下，离线编译器发布如下错误消息以及最外层循环的行号：

Kernel <function> exceeded the Max II. The Kernel's resource usage is estimated to be much larger than FPGA capacity. It performs poorly even if it fits. Reduce resource utilization of the kernel by reducing loop unroll factors within it (if any) or otherwise reduce amount of computation within the kernel.

展开循环并且合并从全局存储器的加载操作使得内核的硬件实现在每个时钟周期执行更多操作。一般而言，您用来提高OpenCL内核性能的方法应达到如下结果：

增加并行操作的数量
增加实现的存储器带宽
增加内核可以在硬件中每个时钟的操作数。

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

您指定完全展开一个具有非常大量迭代的数据依赖循环。因此，您内核的硬件实现可能不符合该FPGA。
您指定完全展开循环，但是该循环边界不是常量。
循环由复杂的控制流组成（例如，包含复杂的数组索引或者在编译时循环包含未知的退出条件）。

对于上述的后面两种情况，离线编译器会发布以下警告：

Full unrolling of the loop is requested but the loop bounds cannot be determined. The loop is not unrolled.

要在这些情况下启用循环展开，请指定#pragma unroll <N> 指令，其中<N>是展开因子。展开因子限制离线编译器展开的迭代数。例如，要防止您内核中的循环被展开，请对该循环添加指令#pragma unroll 1。

请参阅单个Work-Item内核的良好设计实践了解关于构建结构良好的循环的提示。

选择您的语言

使用 Intel.com 搜索

快速链接

最近搜索

高级搜索

仅搜索

Intel® FPGA SDK for OpenCL™ Pro Edition: 最佳实践实践指南

4.2. 展开循环