仅对英特尔可见 — GUID： chn1469549457114
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： chn1469549457114
Ixiasoft
3.3. 局部存储器
局部存储是一个复杂的系统。FPGA不同于具有各级别高速缓存（cache）的典型GPU架构，它通过FPGA内的专用存储器模块实现局部存储器。
局部存储器表征
- Ports：对局部存储器的每次读或写访问都被映射到一个端口。
- Banks：局部存储器的内容可以划分成一个或多个bank，这样局部存储器中的每个bank包含一个数据子集。
- Replicate：Bank由一个或者多个复制（replicate）组成。Bank中的每个复制都具有与其他复制相同的数据。创建复制以有效地支持对局部存储器的多次访问。每个复制都有您的设计可以同时访问的写端口和读端口各一个。如果如果您的局部存储器是double pumped（双泵）式，则每个复制有4个物理端口，其中最多可以有3个是读端口。请参阅Double Pumping（双泵）部分了解更多信息。
- Private copies：复制可以包含一个或多个专用副本（private copies），以允许多个工作组（workgroup）的流水线执行。请参阅局部存储器Bank和专用副本（Private Copies）部分了解更多详细信息。
图 38. 一个或多个M20K块中局部存储器的实现
在您的内核代码中，将局部存储器声明为local类型的变量：
local int lmem[ 1024];英特尔® FPGA SDK for OpenCL™ Offline Compiler可自定义局部存储器的属性，如宽度、深度、bank、复制、专用副本以及互连的数量。连线编译器（offline compiler）会根据您的代码分析访问模式，然后优化局部存储器以最大程度地减少访问争用。
下图说明局部存储器的这些基本属性（大小，宽度，深度，bank，复制和专用副本的数量）：
图 39. 解释局部储存器属性的局部存储器实例
在HTML报告中，将局部存储器的整体状态报告为stall-free（无停顿），stall-free with replication（有复制的无停顿）和potentially inefficient（潜在的效率低下）。
设计高效内核的关键是拥有永无停顿的存储器访问。对于无停顿的存储器配置，则需要保证数据路径中所有可能并发的存储器访问站点都能访问存储器并且无争用。
离线编译器（offline compiler）始终尝试在您的内核代码中为所有局部存储器找出无停顿（stall-free）的配置。然而，在复杂的内核中，离线编译器可能没有足够的信息来推断存储器访问之间是否存在任何冲突。因此，离线编译器会推断局部互连仲裁来调节存储器访问。但是，推断仲裁可能会导致性能下降。请参阅Load-Store Units了解更多信息。
图 40. 复杂的局部存储器系统
离线编译器（offline compiler）并不总是以您指定的确切大小实现局部存储器。由于FPGA RAM块具有特定的维度，offline compiler实现的局部存储器大小向上舍入，达到下一个支持的RAM块维度。请参阅特定器件的信息，来了解关于RAM块的更多信息。
局部存储器Bank和专用副本（Private Copies）
默认情况下，局部存储器banking仅适用于最低维度。具有多个bank可允许实现同时写入。如下代码实例中，循环中的每个局部存储器访问都有一个单独的地址。离线编译器（offline compiler）可以通过推断访问模式，来为lmem创建四个单独的bank。四个单独的bank允许4个同时访问lmem[][]，从而实现无停顿（stall-free）配置。此外，离线编译器为lmem创建两个专用副本（private copies），使得两个工作组的流水线执行同时进行。
#define BANK_SIZE 4
__attribute__((reqd_work_group_size(8, 1, 1)))
kernel void bank_arb_consecutive_multidim (global int* restrict in,
global int* restrict out) {
local int lmem[1024][BANK_SIZE];
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 < BANK_SIZE; i++) {
lmem[((li+i) & 0x7f)][i] = res + i;
res = res >> 1;
}
int rdata = 0;
barrier(CLK_GLOBAL_MEM_FENCE);
#pragma unroll
for (int i = 0; i < BANK_SIZE; i++) {
rdata ^= lmem[((li+i) & 0x7f)][i];
}
out[gi] = rdata;
return;
}下图说明以下局部变量的实现（如Kernel memory Viewer中所示）：
local int lmem[1024][4];
图 41. lmem[1024][4]的实现局部存储器的大小= 32768字节= 2个专用副本（private copies）x (1024 x 4) x 4字节。每个bank的大小是8192字节。
如果专用副本的数量显著增加了您的设计面积，请考虑减少该内核中的barrier（屏障）数量，或者增加max_work_group_size的值来减少推断的专用副本的数量。
您可以使用__attribute__((numbanks(N))指定存储器系统中的bank数。请参阅通过Banking（储存）局部存储器来提高内核性能了解更多信息。
如果您不向存储（bank）在最低维度，则请使用bank_bits属性指定来自存储器地址的位作为所选bank的位。通过使用bank_bits属性，您可以将存储数据分入多个bank，与此同时指定哪个地址位使用所选择的bank。如下实例中，是在第7和第8个位完成存储（banking），而非最低的两个维度上：
#define BANK_SIZE 4
kernel void bank_arb_consecutive_multidim_origin (global int* restrict in,
global int* restrict out) {
local int a[BANK_SIZE][128] __attribute__((bank_bits(8,7),bankwidth(4)));
int gi = get_global_id(0);
int li = get_local_id(0);
int res = in[gi];
#pragma unroll
for (int i = 0; i < BANK_SIZE; i++) {
a[i][((li+i) & 0x7f)] = res + i;
res = res >> 1;
}
int rdata = 0;
barrier(CLK_GLOBAL_MEM_FENCE);
#pragma unroll
for (int i = 0; i < BANK_SIZE; i++) {
rdata ^= a[i][((li+i) & 0x7f)];
}
out[gi] = rdata;
return;
}生成的存储器的视图与第一个示例中的初始视图相同，只是现在存储器的大小是4096 bytes = 2个private copies x (4 x 128) x 4个字节。Kernel Memory Viewer中的Details窗格显示地址位的信息，其中还包含bank_bits信息。
下图显示以下局部变量声明的地址位的信息，如局部存储器报告中所示：
local int a[4][128] __attribute__((bank_bits(8,7),bankwidth(4)));
图 42. 带有需要的bank_bits(8,7)的a[4][128]的地址位信息
Bank-bits的选择可以改变存储器的结构。如果在之前的示例中指定了bank-bits (4,3)，就会导致一个被仲裁的存储器系统。在该banking配置中，局部存储器访问不再以单独的bank为对象。编译器必须构建局部存储器内连来仲裁这些访问，这样就降低了设计性能。
local int a[4][128] __attribute__((bank_bits(4,3),bankwidth(4)));
图 43. 带有需要的bank_bits (4,3)位的a[4][128]局部存储器系统
局部存储器复制（Replication）
要实现一个无停顿的配置，编译器可能需要决定复制一个局部存储器系统来增加可用的读端口数量。对局部存储器系统的每个存储操作与每个复制同时执行，因此每个复制都包含相同的数据。可以从每个复制独立读取。这样就增加了局部存储器系统支持的同时读取操作的数量。
Double Pumping（双泵）
默认情况下，每个局部存储器复制有两个物理端口。而double pumping（双泵）功能允许每个局部存储器复制支持的物理端口最多达到四个。
启动double pumping（双泵）的底层机制是以内核逻辑频率的两倍运行底层M20K。这就使得每个时钟周期可以进行两次读或写操作。从内核逻辑的角度来看，double-pumped存储器有4个有效的物理端口。
图 44. 局部存储器中Double Pumping的硬件架构
通过使能双泵功能，离线编译器在面积与最大频率之间进行权衡。
双泵的优势
- 增加可用物理端口的数量
- 可能通过减少复制来减少RAM的使用
双泵的弊端
- 与单泵（single pumped）配置相比，双泵具有更高的逻辑和延迟
- 可能会降低内核时钟频率
可以使用__attribute__((singlepump))和__attribute__((doublepump))来控制您局部存储器系统的泵配置。请参阅配置局部和专用存储器系统时的内核属性了解更多信息。
如下代码实例说明具有三个读端口和三个写端口的局部存储器的实现。离线编译器启用双泵并复制局部存储器三次以实现stall-free（无停顿）存储器配置。
#define BANK_SIZE 4
kernel void bank_arb_consecutive_multidim_origin (global int* restrict in,
global int* restrict out) {
local int a[BANK_SIZE][128];
int gi = get_global_id(0);
int li = get_local_id(0);
int res = in[gi];
#pragma unroll 1
for (int i = 0; i < BANK_SIZE; i++) {
a[i][li+i] = res + i;
a[gi][li+i] = res + i;
a[gi+i][li] = res + i;
res = res >> 1;
}
int rdata = 0;
barrier(CLK_GLOBAL_MEM_FENCE);
#pragma unroll 1
for (int i = 0; i < BANK_SIZE; i++) {
rdata ^= a[i][li+i];
rdata += a[gi+i][li+i];
rdata += a[gi][li];
}
out[gi] = rdata;
return;
}下图说明如下局部变量声明的实现（如Kernel Memory Viewer中所示）：
local int a[4][128];
图 45. a[4][128]的局部存储器系统局部存储器的大小= 6144字节= 3个复制x 512个字x 4个字节。每个复制都具有相同的存储内容。