仅对英特尔可见 — GUID： hex1566241909270
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： hex1566241909270
Ixiasoft
3.6.2. Load-Store Unit修改程序
编译器会根据内核中的存储器访问模式修改某些LSU。
高速缓存(Cache)
突发合并LSU有时可能包含cache。当存储器访问模式依赖于数据或者看似重复时，就会创建cache。即使需要负载需要相同的数据，cache也不能与其他负载共享。当内核启动时，cache被刷写（flushed），并且比没有cache的等效LSU消耗更多的硬件资源。仅针对非易失性全局指针推断Cache。
kernel void cached (global int * restrict in,
global int * restrict out,
int N) {
int gid = get_global_id(0);
for (int i = 0; i < N; i++) {
out[N*gid + i] = in[i];
}
}
写确认（write-ack）
存在数据依赖项时，突发合并（Burst-coalesced store）LSU有时需要write-acknowledgment（写确认）信号。带有写确认（write-acknowledge）信号的LSU需要额外的硬件资源。如果多个写确认同一存储器，则吞吐量可能会减少。
kernel void write_ack (global int * restrict in,
global int * restrict out,
int N) {
for (int i = 0; i < N; i++) {
if (i < 2)
out[i] = 0; // Burst-coalesced write-ack LSU
out[i] = in[i];
}
} Nonaligned（未对齐）
当突发合并的LSU可以访问未与外部存储器字长对齐的存储器时，将创建一个未对齐的LSU。需要额外的硬件资源来实现未对齐的LSU。如果未对齐的LSU接受太多未对齐的请求，则其吞吐量可能会降低。
kernel void non_aligned (global int * restrict in,
global int * restrict out) {
int i = get_global_id(0);
// Three loads are statically coalesced into one,
// creating a burst-coalesced non-aligned LSU.
int a1 = in[3*i+0];
int a2 = in[3*i+1];
int a3 = in[3*i+2];
// Three stores statically coalesced into one,
// creating a burst-coalesced non-aligned LSU.
out[3*i+0] = a3;
out[3*i+1] = a2;
out[3*i+2] = a1;
}
Never-stall（永无停顿）
如果流水线LSU在没有仲裁的情况下被连接到局部存储器，则会创建永无停顿（never-stall）的LSU，因为对存储器的所有访问都以编译器已知的固定周期数。
__attribute((reqd_work_group_size(1024,1,1)))
kernel void never_stall (global int* restrict in,
global int* restrict out,
int N) {
local int lmem[1024];
int gi = get_global_id(0);
int li = get_local_id(0);
lmem[li] = in[gi]; // Pipelined never-stall LSU
barrier(CLK_GLOBAL_MEM_FENCE);
out[gi] = lmem[li] ^ lmem[li + 1];
}