OpenCL优化方法-Single Work-Item Kernel的良好设计实践

翻译 2017年08月01日 20:01:21
Single Work-Item Kernel的良好设计实践
如果您的OpenCL内核包含循环结构,请遵循英特尔推荐的指导方针,以允许OpenCL Offline Compiler的Intel FPGA SDK可以针对内核进行有效分析。指导离线编译器在循环中执行流水线并行执行时,结构良好的循环尤其重要。

1、避免指针别名混淆使用
在指针参数中尽可能插入restrict关键字。在指针参数中包含restrict关键字可防止离线编译器在非冲突读写操作之间创建不必要的内存依赖关系。考虑一个循环,其中每次迭代从一个数组读取数据,然后将数据写入同一物理内存中的另一个数组。在这些指针参数中没有包含restrict关键字时,离线编译器可能会假定两个数组之间存在依赖关系,从而提取较少的流水线并行性。

2、构建“良好格式”的循环
“良好格式”的循环具有限制与整数相比较的退出条件,并且每次迭代具有一个简单的归纳增量。在内核中运用“良好格式”的循环可以提高性能,因为离线编译器可以有效地分析这些循环。
以下是“良好格式”的循环示例:
for(i=0; i < N; i++)
{
//statements
}
Important:“良好格式”的嵌套循环也有助于最大化内核性能。
以下是“良好格式”的嵌套循环结构示例:
for(i=0; i < N; i++)
{
//statements
for(j=0; j < M; j++)
{
//statements
}
}

3、Minimize Loop-Carried Dependencies(最大限度地减少环路依赖)
下面的循环结构创建一个循环依赖结构,因为每个循环迭代读取由上一次迭代写入的数据。导致的结果就是:每次读取操作都不能继续进行,直到前一次迭代的写操作完成。循环依赖关系的存在降低了离线编译器可以实现的流水线并行性的程度,从而降低了内核性能。
for(int i = 0; i < N; i++)
{
A[i] = A[i - 1] + i;
}
离线编译器对循环执行静态内存依赖性分析,以确定其可以实现的并行性程度。在某些情况下,离线编译器可能会假定两个数组访问之间存在依赖关系,从而提取较少的流水线并行性。如果离线编译器由于未知变量而无法解析编译时的依赖关系,或者数组访问涉及复杂寻址,则会假定循环依赖。
为了尽量减少循环携带的依赖关系,请尽可能遵循以下准则:
  • 避免指针算术。
当内核通过解引用从算术运算导出的指针值来访问数组时,编译器输出次最优方案。
例如,避免以下列方式访问数组:
for(int i = 0; i < N; i++)
{
int t = *(A++);
*A = t;
}
  • 使用简单的数组索引。
避免使用以下复杂类型的数组索引,因为离线编译器无法有效地分析它们,这可能会导致次最优的编译器输出:
-数组索引为非常数变量。
例如,A [K + i],其中i是循环索引变量,K是未知变量。
-多个索引变量在相同的下标位置。
例如,A [i + 2×j],其中i和j是双重嵌套循环的循环索引变量。注意:离线编译器可以有效地分析数组索引A [i] [j],因为索引变量在不同的下标位置。
-非线性索引
例如,A [i&C],其中i是循环索引变量,C是常数或非常数变量。
  • 尽可能在内核中使用具有常量边界的循环。
具有常量边界的循环允许离线编译器有效地执行范围分析。

4、Avoid Complex Loop Exit Conditions(避免复杂的循环退出条件)
离线编译器评估循环退出条件以确定后续循环迭代是否可以进入循环流水线(pipeline)。有时离线编译器需要通过内存访问或复杂操作来评估退出条件。在这些情况下,随后的迭代将不会启动,直到评估完成,会导致整体循环性能的降低。

5、Convert Nested Loops into a Single Loop(将嵌套循环转换为单个循环)
为了最大限度地提高性能,尽可能将嵌套循环合并成单个循环的形式。将嵌套循环重组为单个循环会减少循环迭代之间的硬件占用空间和计算开销。
The following code examples illustrate the conversion of a nested loop into a single loop:
Nested Loop Converted Single Loop
for (i = 0; i < N; i++)                                              {                                                                                //statements                                                          for (j = 0; j < M; j++)                                              {                                                                                  //statements                                                    }                                                                            //statements                                                     } for (i = 0; i < N*M; i++)                                            {                                                                                 //statements                                                    }
6、Declare Variables in the Deepest Scope Possible;(尽可能的在最深的嵌套内声明变量)
(a[N]、b[N]占用资源不同,资源占用情况:a>b)
为了减少实现变量所需的硬件资源,请在循环中使用该变量之前声明该变量。在最深的范围内声明变量可能会最小化数据依赖性和硬件使用率,因为离线编译器(the offline compiler)不需要跨不使用变量的循环保留变量数据。
Consider the following example:
int a[N];
for (int i = 0; i < m; ++i)
{
int b[N];
for (int j = 0; j < n; ++j)
{
// statements
}
}
数组a需要比数组b更多的资源来实现。 为了减少使用的硬件资源,尽可能地将数组声明在内部循环之内,除非需要通过外部循环的迭代维护数据时才将数组声明在内部循环之外。在最深的范围内覆盖变量的所有值也可以减少呈现变量所需的资源。

《OpenCL异构并行计算:原理、机制与优化实践》笔记(二):进入OpenCL的世界(矢量加法)

《OpenCL异构并行计算:原理、机制与优化实践》笔记(二):进入OpenCL的世界(矢量加法)...
  • icamera0
  • icamera0
  • 2017年07月13日 22:46
  • 302

OpenCL中kernel的循环调用

kernel的循环调用主要是涉及缓冲区的创建和主机端命令同步
  • u011028771
  • u011028771
  • 2016年10月09日 10:47
  • 1303

OpenCL 学习step by step (3) 存储kernel文件为二进制

转自:http://www.cnblogs.com/mikewolf2002/archive/2012/09/06/2674125.html 作者:  迈克老狼2012 在教程2中,我们通...
  • lucky_greenegg
  • lucky_greenegg
  • 2013年08月15日 22:34
  • 1617

OpenCL性能优化实例研究系列2:避免Local Memory Bank Conflicts的两个简单方法

转自:http://hi.baidu.com/fsword73/item/51df1fafe6083e268919d39e 作者:  fsword73 Bank Conflicts 是存储...
  • lucky_greenegg
  • lucky_greenegg
  • 2013年08月15日 22:41
  • 1841

C++实战之OpenCL矩阵相乘优化(二)

前言上一篇文章,分析了简单的矩阵相乘在opencl里面的优化kernel代码,每个work-item只负责计算结果矩阵的一个元素。下一步准备每次计算出结果矩阵的块元素,看看计算时间是如何。这个矩阵系列...
  • w401229755
  • w401229755
  • 2017年11月21日 14:51
  • 916

【并行计算-CUDA开发】FPGA 设计者应该学习 OpenCL及爱上OpenCL的十个理由

为什么要学习OpenCL呢?就目前我所从事的医疗超声领域,超声前端的信号处理器一般是通过FPGA或FPGA+DSP来设计的,高端设备用的是FPGA+ GPU架构。传统的设计方法是通过HDL语言来进行设...
  • LG1259156776
  • LG1259156776
  • 2016年09月29日 20:25
  • 2916

关于《OPENCL异构并行计算》中卷积优化的分析

《OPENCL异构并行计算》中讲了如何利用OPENCL进行卷积运算,并给出了使用局部存储器优化的例子,这里对其进行简单分析...
  • qq_20028731
  • qq_20028731
  • 2017年04月16日 17:13
  • 531

Jpeg 库的解码OpenCL优化

libJpeg库解码OpenCL优化这两周在闲暇时基于通用的libjpeg库重新做了一个opencl解码实现。重新熟悉下算法。代码路径https://github.com/jxt1234/platfo...
  • jxt1234and2010
  • jxt1234and2010
  • 2015年04月22日 20:41
  • 2349

opencl学习(四)

由于opencl内核运行时,clEnqueueNDRangeKernel的第5、6个参数global大小和local大小影响计算效率,甚至执行成功情况,想深入了解一下相关的几个参数。参考平台上运行cl...
  • eric41050808
  • eric41050808
  • 2013年08月23日 10:18
  • 7778

OpenCL编程步骤(四):创建内核对象和设置内核参数

内核就是程序中声明的一个函数。对于程序中的任一函数,都可以通过加上限定符__kernel将其标识为内核。内核对象中封装了程序中的某个__kernel函数以及执行此函数时所需的参数。 ...
  • u013684730
  • u013684730
  • 2015年05月27日 16:29
  • 2145
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:OpenCL优化方法-Single Work-Item Kernel的良好设计实践
举报原因:
原因补充:

(最多只允许输入30个字)