仅对英特尔可见 — GUID: ewa1397680164688
Ixiasoft
产品终止通知
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版最佳实践指南修订历史
仅对英特尔可见 — GUID: ewa1397680164688
Ixiasoft
9.3. 存储器访问考量
Intel® 建议采用可以提高存储器访问效率并减少OpenCL™内核面积使用的内核编程策略。
- 最小化对外部存储器访问点的数量。
如果可能,构建您的内核,使其从一个位置读取其输入,在内部处理数据,然后将输出写入另一个位置。
- 不要依赖全局存储器访问,而是尽可能将内核构建未具有移位寄存器推断的单个work-item。
- 与其创建将数据写入外部存储器的内核和从外部存储器读取数据的内核,不如实现内核之间的 Intel® FPGA SDK for OpenCL™ 通带扩展,用于直接数据传输。
- 如果您的OpenCL应用程序包含许多单独的常量访问,请使用__constant而不是__global const声明相应的指针。使用__global const的声明为每个负载或存储操作创建一个专用cache。但是,使用__constant的声明只会在片上创建一个常量cache。
警告:如果您的内核以Cyclone® V器件(例如,Cyclone V SoC)为目标,声明__constant指针内核自变量可能降低FPGA的性能。
- 如果您的内核传递少量常数自变量,请将它们作为值而非指向全局存储器的指针传递。
例如,不传递__constant int * coef,但是解除对索引0到10的coef引用,并将coef作为value(int16 coef)传递。如果coef 是唯一的__constant指针自变量,将其作为值传递会完全消除常量cache以及相应的负载和存储操作。
- 在流水线循环内有条件地shifting(移位)大型移位寄存器会导致创建低效的硬件。例如,if (K > 5)条件存在时,以下内核会消耗更多资源:
#define SHIFT_REG_LEN 1024 __kernel void bad_shift_reg (__global int * restrict src, __global int * restrict dst, int K) { float shift_reg[SHIFT_REG_LEN]; int sum = 0; for (unsigned i = 0; i < K; i++) { sum += shift_reg[0]; shift_reg[SHIFT_REG_LEN-1] = src[i]; // This condition will cause sever area bloat. if (K > 5) { #pragma unroll for (int m = 0; m < SHIFT_REG_LEN-1 ; m++) { shift_reg[m] = shift_reg[m + 1]; } } dst[i] = sum; } }
注意: 有条件地accessing(访问)移位寄存器不会降低硬件效率。如果有必要在您的内核中实现较大移位寄存器的有条件移位。