仅对英特尔可见 — GUID: rfr1469543500580
Ixiasoft
仅对英特尔可见 — GUID: rfr1469543500580
Ixiasoft
3.4.3. 嵌套循环
由于循环迭代依顺序进行, Intel® FPGA SDK for OpenCL™ Offline Compiler不推断流水线执行。因此,相对于下一个内部循环,外部循环(outer loop)迭代可能是无序的,因为内部循环迭代的次数可能针对各种外部循环迭代而不同。
要解决无序的外循环迭代问题,请将内部循环的下界(lower bound)和上界(upper bound) 设计为外部循环迭代之间不发生改变。
单个Work-Item执行
为确保FPGA上高吞吐量单个基于work-item的kernel执行, Intel® FPGA SDK for OpenCL™ Offline Compiler必须在任何给定时间并行处理多个流水线阶段。该并行性(parallelism)是通过流水线化循环的迭代来实现。
请参考如下简单实例代码,它显示为单个work-item的累加:
1 kernel void accum_swg (global int* a,
global int* c,
int size,
int k_size) {
2 int sum[1024];
3 for (int k = 0; k < k_size; ++k) {
4 for (int i = 0; i < size; ++i) {
5 int j = k * size + i;
6 sum[k] += a[j];
7 }
8 }
9 for (int k = 0; k < k_size; ++k) {
10 c[k] = sum[k];
11 }
12 }
下图说明i的每次迭代如何进入块:
您在观察外部循环时会看到,II值为1也意味着线程的每次迭代可在每个时钟周期进入。在本实例中,认为k_size为20和size为4。这样对于首8个时钟周期来说是正确的,因为外部循环迭代0到7可以进入,而且无任何下游(downstream)中止它。一旦 线程0进入内部循环,它需要4次迭代才能完成。线程1到8无法进入内部循环,并且它们会被线程0停顿4个周期。线程1在线程0的迭代完成后才进入内部循环。因此, 线程9在时钟周期13进入外循环。线程9到20每四个时钟周期进入循环, 这个由size的值决定。通过该实例,您可以观察到外循环的动态启动间隔大于静态预测的启动间隔1,并且它是内循环的行程计数的函数。
非线性执行
循环结构不支持线性执行。以下代码实例显示外部循环i包含两个发散(divergent)的内部循环。外部循环的每次迭代都可能执行一个内部循环或者另一个内部循环,这是一种非线性执行模式。
__kernel void structure (__global unsigned* restrict output1,
__global unsigned* restrict output2,
int N) {
for (unsigned i = 0; i < N; i++) {
if ((i & 3) == 0) {
for (unsigned j = 0; j < N; j++) {
output1[i+j] = i * j;
}
}
else
{
for (unsigned j = 0; j < N; j++) {
output2[i+j] = i * j;
}
}
}
}
Serial Regions(串行区域)
当内部循环访问导致外部循环依赖时,嵌套循环中就可能出现串行区域。由于数据或者存储器依赖性,该内部循环就成为外部循环迭代中的串行区域。
稳定状态下,外部循环的II = 内部循环的II *内部循环的形成计数。对于II大于1的内部循环,以及无串行执行区域的外部循环,就可能从外部循环交错线程。
请参考如下代码实例:
kernel void serially_execute (global int * restrict A,
global int * restrict B,
global int * restrict result,
unsigned N) {
int sum = 0;
for (unsigned i = 0; i < N; i++) {
int res;
for (int j = 0; j < N; j++) {
sum += A[i*N+j];
}
sum += B[i];
}
*result = sum;
}