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

下载 PDF
ID 683521
日期 9/26/2022
Public
文档目录
文档目录 x
产品终止通知 1. Intel® FPGA SDK for OpenCL™ Pro Edition最佳实践指南介绍 2. 查看您Kernel的report.html文件 3. OpenCL内核设计概念 4. OpenCL内核设计最佳实践 5. 分析（Profiling）您的内核来识别性能瓶颈 6. 提高单个Work-Item内核性能的策略 7. 提高NDRange内核数据处理效率的策略 8. 提高存储器访问效率的策略 9. 优化FPGA面积使用的策略 10. 优化英特尔 Stratix 10 OpenCL设计的策略 11. 提高主机应用程序性能的策略 12. Intel® FPGA SDK for OpenCL™ Pro版最佳实践指南存档 A. Intel® FPGA SDK for OpenCL™ Pro版最佳实践指南修订历史
1. Intel® FPGA SDK for OpenCL™ Pro Edition最佳实践指南介绍 x
1.1. FPGA概述 1.2. 流水线 1.3. 单个work-item Kernel与NDRange Kernel
2. 查看您Kernel的report.html文件 x
2.1. 高层次设计报告布局 2.2. 查看Summary Report 2.3. 查看设计中吞吐量瓶颈 2.4. 使用视图 2.5. 分析吞吐量 2.6. 查看区域信息 2.7. 基于HTML Report中的信息优化OpenCL设计实例 2.8. 访问JSON格式的HLD FPGA报告
2.4. 使用视图 x
2.4.1. System Viewer的功能 2.4.2. Kernel Memory Viewer的功能 2.4.3. 调度Viewer的功能
2.4.1. System Viewer的功能 x
2.4.1.1. 查看系统信息 2.4.1.2. 查看全局存储器信息 2.4.1.3. 查看块信息 2.4.1.4. 查看集群信息
2.5. 分析吞吐量 x
2.5.1. 查看循环信息
2.6. 查看区域信息 x
2.6.1. 板级接口的区域报告消息 2.6.2. 针对功能开销的区域报告消息 2.6.3. 针对状态（State）的区域报告消息 2.6.4. 区域报告消息中的反馈 2.6.5. Area Report Message中的Private Variable Storage
2.6.5. Area Report Message中的Private Variable Storage x
2.6.5.1. 针对Constant Memory的Area Report Message
3. OpenCL内核设计概念 x
3.1. 内核 3.2. 全局存储互连 3.3. 局部存储器 3.4. 单个Work-Item内核中的循环 3.5. 通道 3.6. Load-Store Units
3.3. 局部存储器 x
3.3.1. 更改存储器访问模式的实例
3.4. 单个Work-Item内核中的循环 x
3.4.1. 启动间隔与最大频率之间的权衡 3.4.2. 循环携带的依赖项影响了循环的启动间隔 3.4.3. 嵌套循环 3.4.4. 循环推测 3.4.5. 循环融合（Loop Fusion） 3.4.6. 循环瓶颈（Loop Bottleneck）
3.4.3. 嵌套循环 x
3.4.3.1. 使用loop_coalesce减少嵌套循环占用的面积
3.6. Load-Store Units x
3.6.1. Load-Store Unit类型 3.6.2. Load-Store Unit修改程序 3.6.3. 控制Load-Store Units 3.6.4. 何时使用各种LSU
4. OpenCL内核设计最佳实践 x
4.1. 通过 Intel® FPGA SDK for OpenCL™ 通道或OpenCL管道传输数据 4.2. 展开循环 4.3. 优化浮点操作 4.4. 分配对齐的存储器 4.5. 使用/不使用填充来对齐结构体 4.6. 矢量类型单元保持相似的结构 4.7. 避免指针别名 4.8. 避免昂贵的函数 4.9. 避免Work-Item ID依赖的向后分支
4.1. 通过 Intel® FPGA SDK for OpenCL™ 通道或OpenCL管道传输数据 x
4.1.1. 通道和管道的表征 默认行为 多个OpenCL内核的并发执行 隐式内核同步 调用之中的数据持久性 强制实施的Work-Item顺序 4.1.2. 通道和管道的执行顺序 4.1.3. 优化通道或管道的缓冲推断 4.1.4. 通道和管道的最佳实践
4.3. 优化浮点操作 x
4.3.1. 浮点与定点表示
5. 分析（Profiling）您的内核来识别性能瓶颈 x
5.1. 分析您内核的最佳实践 5.2. 使用Performance Counters (-profile)设施注入（instrument）Kernel Pipeline 5.3. 获取运行期间分析数据 5.4. 减少分析时的面积资源使用 5.5. 时间（Temporal）性能集合 5.6. 性能数据类型 5.7. 解释分析信息 5.8. 示例OpenCL设计场景的Profiler分析 5.9. Intel® FPGA Dynamic Profiler for OpenCL™ 限制
5.3. 获取运行期间分析数据 x
5.3.1. 调用Profiler Runtime Wrapper 5.3.2. 使用Intel® VTune™ Profiler查看分析数据
5.3.1. 调用Profiler Runtime Wrapper x
5.3.1.1. 拆分执行和数据后处理
5.5. 时间（Temporal）性能集合 x
5.5.1. 分析（Profiling）Autorun内核
5.7. 解释分析信息 x
5.7.1. 停顿、占用率、带宽 5.7.2. 停顿通道 5.7.3. 通道深度
5.8. 示例OpenCL设计场景的Profiler分析 x
5.8.1. 高停顿百分比 5.8.2. 较低的占用量百分比 5.8.3. 高停顿和高占用量百分比 5.8.4. 无停顿、低占用量百分比和低带宽 5.8.5. 无停顿、高占用量百分比和低带宽 5.8.6. 高停顿和低占用量百分比
6. 提高单个Work-Item内核性能的策略 x
6.1. 根据优化报告反馈解决单个Work-Item内核依赖项 6.2. 单个Work-Item内核的良好设计实践
6.1. 根据优化报告反馈解决单个Work-Item内核依赖项 x
6.1.1. 删除循环携带的依赖项 6.1.2. 松弛循环携带的依赖性 6.1.3. 将循环携带的依赖项转移到局部存储器 6.1.4. 通过推断移位寄存器来松弛循环携带的依赖项 6.1.5. 移除由于对存储器阵列的访问而引起的循环依赖
7. 提高NDRange内核数据处理效率的策略 x
7.1. 指定最大工作组大小或者需要的工作组大小 7.2. 内核矢量化 7.3. 多个计算单元 7.4. 计算单元复制与内核SIMD矢量化合并 7.5. 在HTML报告中查看内核属性和循环展开状态
7.3. 多个计算单元 x
7.3.1. 计算单元复制与内核SIMD向量化
8. 提高存储器访问效率的策略 x
8.1. 优化存储器访问的常规指导 8.2. 优化全局存储器访问 8.3. 使用常量、局部或专用存储器执行内核计算 8.4. 通过Banking（储存）局部存储器来提高内核性能 8.5. 通过控制存储器赋值因子来优化对局部存储器的访问 8.6. 最小化循环流水线的存储器依赖 8.7. 静态存储器合并
8.2. 优化全局存储器访问 x
8.2.1. Contiguous Memory Access（连续存储访问） 8.2.2. 全局存储器手动分区 8.2.3. 优化全局存储器的一个或多个Bank
8.2.2. 全局存储器手动分区 x
8.2.2.1. 异构存储器缓冲区
8.3. 使用常量、局部或专用存储器执行内核计算 x
8.3.1. 常量缓存存储器 8.3.2. 将数据预加载到局部存储器 8.3.3. 在专用存储器中存储可变量和数组
8.4. 通过Banking（储存）局部存储器来提高内核性能 x
8.4.1. 基于数组索引优化局部存储器Bank的几何配置
9. 优化FPGA面积使用的策略 x
9.1. 编译考量 9.2. 电路板变体选择考量 9.3. 存储器访问考量 9.4. 算法操作考量 9.5. 数据类型选择考量
10. 优化英特尔 Stratix 10 OpenCL设计的策略 x
10.1. 减少通道开销 10.2. 优化循环控制 10.3. 简化对局部存储器的存储访问 10.4. 重复使用数据的片上存储 10.5. 优化数据路径控制 10.6. 创建RTL模块
10.1. 减少通道开销 x
10.1.1. 减少内核数量 10.1.2. 使用单个内核来描述脉动阵列 10.1.3. 使用非阻塞通道（Non-Blocking）通道
10.2. 优化循环控制 x
10.2.1. 简化英特尔 Stratix 10 OpenCL设计中的循环携带依赖性
10.6. 创建RTL模块 x
10.6.1. 复位建议
11. 提高主机应用程序性能的策略 x
11.1. 多线程主机应用程序 11.2. 使用硬件内核调用队列
11.2. 使用硬件内核调用队列 x
11.2.1. 双缓冲主机应用程序利用内核调用队列
产品终止通知
1. Intel® FPGA SDK for OpenCL™ Pro Edition最佳实践指南介绍
2. 查看您Kernel的report.html文件
3. OpenCL内核设计概念
4. OpenCL内核设计最佳实践
5. 分析（Profiling）您的内核来识别性能瓶颈
6. 提高单个Work-Item内核性能的策略
7. 提高NDRange内核数据处理效率的策略
8. 提高存储器访问效率的策略
9. 优化FPGA面积使用的策略
10. 优化英特尔 Stratix 10 OpenCL设计的策略
11. 提高主机应用程序性能的策略
12. Intel® FPGA SDK for OpenCL™ Pro版最佳实践指南存档
A. Intel® FPGA SDK for OpenCL™ Pro版最佳实践指南修订历史

4.1.1. 通道和管道的表征

要在您的OpenCL™内核程序中实现通道或管道，请记住它们各自特定于 Intel® FPGA SDK for OpenCL™ 的表征。

默认行为

通道的默认行为是阻塞（blocking）。管道的默认行为是非阻塞（nonblocking）。

多个OpenCL内核的并发执行

您可并发执行多个OpenCL内核。要使能并发执行，请修改主机代码以实例化多个命令队列。每个并发执行内核都与一个单独的命令队列相关联。

重要：

特定于管道的考量：

Intel® FPGA SDK for OpenCL™ 编程指南的确保与其它OpenCL SDK的兼容性中概述了OpenCL管道修改的内容。OpenCL管道修改允许您在SDK上运行您的内核。但不会最大化内核的吞吐量。OpenCL Specification版本2.0要求管道写入在先，然后才进行管道读取，这样该内核就不会从空的管道读取。因此，内核不能并发执行。然而因为 Intel® FPGA SDK for OpenCL™ 支持并发执行，所以可以修改您的主机应用程序和内核程序来利用该功能。修改后可以增加应用程序的吞吐量；但是您不能再将内核转接（port）到另外一个SDK。尽管存在此限制，但修改很少，并且也不需要太大的工作量来维护这两类代码。

要启用包含管道的内核的并发执行，请用blocking属性（即，__attribute__((blocking))）取代您内核代码中的depth属性。blocking属性会在read_pipe和write_pipe函数调用中引入阻塞（blocking）行为。该调用站点阻塞（block）内核执行，一直到管道的另一端准备就绪。

如果将blocking属性和depth属性都添加到您的内核中，当管道为空时，read_pipe仅调用一个块；而当管道填满时，write_pipe仅调用一个块。阻塞行为（Blocking behavior）会导致内核之间的隐式同步， 这样就迫使内核彼此之间以锁步（lock step）方式运行。

隐式内核同步

通过阻塞通道调用或阻塞管道调用隐式同步内核。请考虑如下示例：

表 14.  阻塞通道和管道调用实现内核同步
具有阻塞通道调用的内核 具有阻塞管道调用的内核 
channel int c0; 

__kernel
void producer (__global int * in_buf) 
{ 
  for (int i = 0; i < 10; i++) 
  {   
    write_channel_intel (c0, in_buf[i]);
  } 
}

__kernel
void consumer (__global int * ret_buf) 
{                 
  for (int i = 0; i < 10; i++) 
  { 
    ret_buf[i] = read_channel_intel(c0);
  }
} 
__kernel
void producer (__global int * in_buf,
  write_only pipe int __attribute__
  ((blocking)) c0) 
{
  for (int i = 0; i < 10; i++) 
{
    write_pipe (c0, &in_buf[i]);
  }
}

__kernel
void consumer (__global int * ret_buf,
  read_only pipe int __attribute__
  ((blocking)) c0) 
{
  for (int i = 0; i < 10; i++) 
  {
    int x;
    read_pipe (c0, &x);
    ret_buf[i] = x;
  }
}

您可以同步内核，以便在每次循环迭代期间，producer内核写入数据，而consumer内核读取数据。如果producer中的write_channel_intel或者write_pipe调用没有写入任何数据，consumer阻塞并等待在read_channel_intel或read_pipe调用，直到 producer发送有效数据，反之亦然。

调用之中的数据持久性

在write_channel_intel调用将数据写入通道，或者write_pipe调用将数据写入管道后，数据将在工作组和NDRange调用中保持不变。work-item写入通道或管道的数据保留在该通道或管道中，直到另外一个work-item从中读取。此外，通道或管道中的数据顺序等同于对该通道或管道的写入操作的顺序，然而该顺序与执行写操作的work-item无关。

例如，如果多个work-item尝试同时访问通道或管道，则只有一个work-item可以访问。write_channel_intel调用或write_pipe调用将名为DATAX的特定work-item数据分别写入通道或管道。与此类似的是，第一个访问该通道或管道的work-item从中读取DATAX。这种顺序排列的写操作使得通道和管道成为内核之间共享数据的有效方式。

强制实施的Work-Item顺序

SDK强制实施work-item顺序来保持通道或管道的读写操作一致性。

相关信息
二级标题