仅对英特尔可见 — GUID： ase1566243064294
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： ase1566243064294
Ixiasoft
3.6.1. Load-Store Unit类型
编译器可以根据推断的存储器访问模式（pattern），目标平台上可用的存储器类型，以及是对局部还是全局存储器的存储器访问生成几种不同类型的load-store units (LSUs) 。 Intel® FPGA SDK for OpenCL™ Offline Compiler可以生成如下类型的LSU：
- Burst-Coalesced（突发合并） Load-Store Units
- 预取Load-Store Units
- 流水线Load-Store Units
- 恒定流水线Load-Store Units
- 原子流水线（Atomic-Pipelined）Load-Store Units
Burst-Coalesced（突发合并） Load-Store Units
Burst-coalesced（突发合并）LSU是为了访问全局存储器，由编译器例化的默认LSU类型。将请求缓冲直到可形成最大型的可能突发。Burst-coalesced LSU可提供对全局存储器的高效访问，但是它需要大量的FPGA资源。
kernel void burst_coalesced (global int * restrict in,
global int * restrict out) {
int i = get_global_id(0);
int value = in[i/2]; // Burst-coalesced LSU
out[i] = value;
}
取决于存储器访问模型和其他属性，编译器可能会以如下方式修改burst-coalesced LSU：
预取Load-Store Units
预取LSU例化FIFO，该FIFO从存储器突发读取大型块，从而保持FIFO充满有效数据，这些有效数据是基于之前的地址并认为是连续读取的。支持Non-contiguous（非连续）读取，但是会导致刷新和重新填充FIFO的代价。预取LSU仅适用于非易失性全局指针。
kernel void prefetching (global int * restrict in,
global int * restrict out,
int N) {
int res = 1;
for (int i = 0; i < N; i++) {
int v = in[i]; // Prefetching LSU
res ^= v;
}
out[0] = res;
}流水线Load-Store Units
流水线LSU用于访问局部存储器。它们会立即提交收到的请求。存储器访问被流水线化，因此一次可以运行多个请求。如果LSU和局部存储器之间没有仲裁，则创建流水线永不停顿的LSU。
__attribute((reqd_work_group_size(1024,1,1)))
kernel void local_pipelined (global int* restrict in,
global int* restrict out) {
local int lmem[1024];
int gi = get_global_id(0);
int li = get_local_id(0);
int res = in[gi];
for (int i = 0; i < 4; i++) {
lmem[li - i] = res; // pipelined LSU
res >>= 1;
}
barrier(CLK_GLOBAL_MEM_FENCE);
res = 0;
for (int i = 0; i < 4; i++) {
res ^= lmem[li - i]; // pipelined LSU
}
out[gi] = res;
}
编译器可能以如下方式修改局部流水线化（local-pipelined）：
编译器还可以为全局存储器访问推断一个流水线LSU，但已经证明这些全局存储器访问并不频繁。编译器使用流水线LSU进行此类访问，因为流水线LSU比其它LSU类型小。虽然流水线LSU的吞吐量可能较低，但该吞吐量权衡还是可以接受的，因为这些存储器访问并不频繁。
kernel void global_infrequent (global int * restrict in,
global int * restrict out,
int N) {
int a = 0;
if (get_global_id(0) == 0)
a = in[0]; // Pipelined LSU
for (int i = 0; i < N; i++) {
out[i] = in[i] + a;
}
}
恒定流水线Load-Store Units
恒定流水线（constant-pipelined）LSU是主要用于从恒定缓存中进行读取的流水线LSU。恒定流水线LSU比突发合并（burst-coalesced）LSU占用更少的面积。恒定流水线LSU很大程度上取决于读取是否命中恒定缓存。缓存不命中是很昂贵的。
kernel void constant_pipelined (constant int *src,
global int *dst) {
int i = get_global_id(0);
dst[i] = src[i]; // Constant pipelined LSU
}有关恒定缓存的信息，请参阅常量缓存存储器。
原子流水线（Atomic-Pipelined）Load-Store Units
原子流水线LSU被用于全部原子操作。使用原子操作可以显著降低内核性能。
kernel void atomic_pipelined (global int* restrict out) {
atomic_add(&out[0], 1); // Atomic LSU
}