仅对英特尔可见 — GUID: ahd1504718313058
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: ahd1504718313058
Ixiasoft
3.3.1. 更改存储器访问模式的实例
以下是一个简单OpenCL内核的示例代码:
kernel void big_lmem_4r_4w_nosplit (global int* restrict in,
global int* restrict out) {
local int lmem[4][1024];
int gi = get_global_id(0);
int gs = get_global_size(0);
int li = get_local_id(0);
int ls = get_local_size(0);
int res = in[gi];
#pragma unroll
for (int i = 0; i < 4; i++) {
lmem[i][(li*i) % ls] = res;
res >>= 1; }
// Global memory barrier
barrier(CLK_GLOBAL_MEM_FENCE);
res = 0;
#pragma unroll
for (int i = 0; i < 4; i++) {
res ^= lmem[i][((ls-li)*i) % ls]; }
out[gi] = res;
}
在System Viewer报告中,该示例的系统视图中突出显示了可停顿的加载和存储。
图 46. 本示例的系统视图
图 47. 本示例的面积报告
图 48. 本示例的内核存储器视图
可观察到只创建了两个存储器bank,并在加载和存储操作之间对第一个bank进行了高度仲裁。现在,切换banking指示到第二个维度,如以下示例代码中所示:
kernel void big_lmem_4r_4w_nosplit (global int* restrict in,
global int* restrict out) {
local int lmem[1024][4];
int gi = get_global_id(0);
int gs = get_global_size(0);
int li = get_local_id(0);
int ls = get_local_size(0);
int res = in[gi];
#pragma unroll
for (int i = 0; i < 4; i++) {
lmem[(li*i) % ls][i] = res;
res >>= 1;
}
// Global memory barrier
barrier(CLK_GLOBAL_MEM_FENCE);
res = 0;
#pragma unroll
for (int i = 0; i < 4; i++) {
res ^= lmem[((ls-li)*i) % ls][i];
}
out[gi] = res;
}
在内核存储器查看器中,您可以观察到现在创建了4个存储器bank,具有单独的加载存储单元。所有加载存储指令都是无停顿的。
图 49. 更改Banking指示后的内核存储器查看器实例
图 50. 更改Banking指示后的面积报告的实例