Intel OpenCL 之 Pipeline(四):Pipeline不理想的几种情况
pipeline不理想的情况主要有两类,一类是影响II
的,一类是不影响II
的。影响II
的会导致II
值大于1,不影响II
的称为Serial Regions
。
通常来说,我们必须保证 critical loop
的II
为1,而非 critical loop
,比如外层循环,在某些条件下可以适当放低II要求。
pipeline不理想通常是由loop-carried dependency
导致,因此本文中先介绍loop-carried dependency
,再介绍两类pipeline不理想的情况。这里只做简单介绍,具体优化后面再详细说明。
loop-carried dependency
Loop-carried dependency可分为两类:
- data dependency
- memory dependency
数据依赖:
下面的例子中,变量sum
存在数据依赖,导致serial region
。
这里,serial region
即内层循环体。编译器为了保证代码运行的正确性,会在新一次迭代进入内层循环体之前,强制保证当前迭代内层循环体已经执行完成,即内层循环完成对变量sum
的更新。
kernel void serially_execute (global int * restrict A,
global int * restrict B,
global int * restrict result, unsigned N)
{
int sum = 0;
for (unsigned i = 0; i < N; i++) {
int res;
for (int j = 0; j < N; j++) {
sum += A[i*N+j];
}
sum += B[i];
}
*result = sum;
}
Loop Report:
-+ Loop "Block1" (file k.cl line 9)
| Pipelined with successive iterations launched every 2 cycles due to:
|
| Pipeline structure: every terminating loop with subloops has iterations launched at least 2 cycles apart.
| Having successive iterations launched every two cycles should still lead to good performance if the inner loop is pipelined well and has sufficiently high number of iterations.
|
| Iterations executed serially across the region listed below.
| Only a single loop iteration will execute inside the listed region.
| This will cause performance degradation unless the region is pipelined well (can process an iteration every cycle).
|
| Loop "Block2" (file k.cl line 10)
| due to:
| Data dependency on variable sum (file k.cl line 7)
|
|
|-+ Loop "Block2" (file k.cl line 10)
Pipelined well. Successive iterations are launched every cycle.
访存依赖:
下面的例子中,由于对global memory
的访问存在依赖关系,导致II值大于1。
我们知道RAM只有一个读端口和一个写端口,只能同时进行一次读操作和一次写操作。当然,Intel FPGA可以构造有3个读端口,一个写端口的RAM,但是也不能像寄存器一样,做到对任意位置的任意读写。编译器为了保证以RAM存储的变量的正确性,对同一变量(变量名),若在一个for循环内既有加载操作,又有存储操作,则会强制保证其先后顺序,因此才会存在仿存依赖的问题。
此外,我们也应该了解到,对global memory
的访问会产生很大延时。
#define N 128
__kernel void unoptimized( __global int* restrict A )
{
for(unsigned i = 0; i < N; i++)
A[N-i] = A[i];
}
Loop Report:
-+ Loop "Block1" (file unoptimized4.cl line 5)
| Pipelined with successive iterations launched every 324 cycles due to:
|
| Memory dependency on Load Operation from: (file unoptimized4.cl line 6)
| Store Operation (file unoptimized4.cl line 6)
| Largest Critical Path Contributors:
| 49%: Load Operation (file unoptimized4.cl line 6)
| 49%: Store Operation (file unoptimized4.cl line 6)
Initiation Interval(II)大于1
理想情况下,II应为1,即优化报告中给出:
Pipelined well. Successive iterations are launched every cycle.
当II大于1时,优化报告会显示以下信息:
Successive iterations launched every <N> cycles due to:
其中,due to后面会给出loop-carried dependency的位置。
Serial Regions
当存在serial regions时,优化报告会给出以下信息:
Loop Report:
-+ Loop "Block1" (file k.cl line 9)
| Pipelined with successive iterations launched every cycles:
|
| Iterations executed serially across the region listed below.
| Only a single loop iteration will execute inside the listed region.
| This will cause performance degradation unless the region is pipelined well (can process an iteration every cycle).
|
| Loop "Block2" (file k.cl line 10)
| due to:
| Data dependency on variable sum (file k.cl line 7)
serial region导致loop的执行成下图的状态:
pipeline-41
而正常能pipeline良好的loop执行是这样的:
pipeline-42
可见,如果内层循环的迭代次数N足够大(远大于内层循环的latency),那么serial regions对性能的影响也不会太大。
参考
[Intel FPGA SDK for OpenCL Best Practices Guide]
- JavaScript 教程
- JavaScript 编辑工具
- JavaScript 与HTML
- JavaScript 与Java
- JavaScript 数据结构
- JavaScript 基本数据类型
- JavaScript 特殊数据类型
- JavaScript 运算符
- JavaScript typeof 运算符
- JavaScript 表达式
- JavaScript 类型转换
- JavaScript 基本语法
- JavaScript 注释
- Javascript 基本处理流程
- Javascript 选择结构
- Javascript if 语句
- Javascript if 语句的嵌套
- Javascript switch 语句
- Javascript 循环结构
- Javascript 循环结构实例
- Javascript 跳转语句
- Javascript 控制语句总结
- Javascript 函数介绍
- Javascript 函数的定义
- Javascript 函数调用
- Javascript 几种特殊的函数
- JavaScript 内置函数简介
- Javascript eval() 函数
- Javascript isFinite() 函数
- Javascript isNaN() 函数
- parseInt() 与 parseFloat()
- escape() 与 unescape()
- Javascript 字符串介绍
- Javascript length属性
- javascript 字符串函数
- Javascript 日期对象简介
- Javascript 日期对象用途
- Date 对象属性和方法
- Javascript 数组是什么
- Javascript 创建数组
- Javascript 数组赋值与取值
- Javascript 数组属性和方法
- 超性感的React Hooks(五):自定义hooks
- 三、变量对象
- Java编程 经验技巧汇总
- 关于IP地址的一些相关知识点
- vmware的三种网络模式
- 我的 Chrome 版本不支持生成二维码,30 分钟怒怼了一个插件,附源码
- Android开发(第一行代码 第二版) 常见异常和解决办法(基于Android Studio)(二)
- 大厂Java项目如何进行Maven多模块管理
- Android开发 经验技巧汇总(基于Android Studio)(一)
- 老板逼我用 Git,本地指令介绍
- Python全栈(三)数据库优化之5.MySQL自关联、外键与Python操作MySQL
- Android开发 经验技巧汇总(基于Android Studio)(二)
- 表格滑动和图片链接,mdnice安排上了!
- 数据库编程 MySQL 常见异常和解决办法
- 持续集成利器,GitHub Actions