Intel OpenCL 之 Pipeline（四）：Pipeline不理想的几种情况

AI异构

发布于 2020-07-29 11:09:45

1.1K0

发布于 2020-07-29 11:09:45

文章被收录于专栏：AI异构 AI异构AI异构

pipeline不理想的情况主要有两类，一类是影响II的，一类是不影响II的。影响II的会导致II值大于1，不影响II的称为Serial Regions。

通常来说，我们必须保证 critical loop 的II为1，而非 critical loop，比如外层循环，在某些条件下可以适当放低II要求。

pipeline不理想通常是由loop-carried dependency导致，因此本文中先介绍loop-carried dependency，再介绍两类pipeline不理想的情况。这里只做简单介绍，具体优化后面再详细说明。

loop-carried dependency

Loop-carried dependency可分为两类：

data dependency
memory dependency

数据依赖：

下面的例子中，变量sum存在数据依赖，导致serial region。

这里，serial region即内层循环体。编译器为了保证代码运行的正确性，会在新一次迭代进入内层循环体之前，强制保证当前迭代内层循环体已经执行完成，即内层循环完成对变量sum的更新。

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;
}

Loop Report:
-+ Loop "Block1" (file k.cl line 9)
| Pipelined with successive iterations launched every 2 cycles due to:
|
|     Pipeline structure: every terminating loop with subloops has iterations launched at least 2 cycles apart.
|     Having successive iterations launched every two cycles should still lead to good performance if the inner loop is pipelined well and has sufficiently high number of iterations.
|
| Iterations executed serially across the region listed below.
| Only a single loop iteration will execute inside the listed region.
| This will cause performance degradation unless the region is pipelined well (can process an iteration every cycle).
|
|     Loop "Block2" (file k.cl line 10)
|     due to:
|     Data dependency on variable sum (file k.cl line 7)
|
|
|-+ Loop "Block2" (file k.cl line 10)
    Pipelined well. Successive iterations are launched every cycle.

访存依赖：

下面的例子中，由于对global memory的访问存在依赖关系，导致II值大于1。

我们知道RAM只有一个读端口和一个写端口，只能同时进行一次读操作和一次写操作。当然，Intel FPGA可以构造有3个读端口，一个写端口的RAM，但是也不能像寄存器一样，做到对任意位置的任意读写。编译器为了保证以RAM存储的变量的正确性，对同一变量（变量名），若在一个for循环内既有加载操作，又有存储操作，则会强制保证其先后顺序，因此才会存在仿存依赖的问题。

此外，我们也应该了解到，对global memory的访问会产生很大延时。

#define N 128

__kernel void unoptimized( __global int* restrict A )
{
  for(unsigned i = 0; i < N; i++)
    A[N-i] = A[i];
}

Loop Report:
-+ Loop "Block1" (file unoptimized4.cl line 5)
 |  Pipelined with successive iterations launched every 324 cycles due to:
 |
 |  Memory dependency on Load Operation from: (file unoptimized4.cl line 6)
 |     Store Operation (file unoptimized4.cl line 6)
 |  Largest Critical Path Contributors:
 |     49%: Load Operation (file unoptimized4.cl line 6)
 |     49%: Store Operation (file unoptimized4.cl line 6)

Initiation Interval（II）大于1

理想情况下，II应为1，即优化报告中给出：

Pipelined well. Successive iterations are launched every cycle.

当II大于1时，优化报告会显示以下信息：

Successive iterations launched every <N> cycles due to:

其中，due to后面会给出loop-carried dependency的位置。

Serial Regions

当存在serial regions时，优化报告会给出以下信息：

Loop Report:
-+ Loop "Block1" (file k.cl line 9)
| Pipelined with successive iterations launched every cycles:
|
| Iterations executed serially across the region listed below.
| Only a single loop iteration will execute inside the listed region.
| This will cause performance degradation unless the region is pipelined well (can process an iteration every cycle).
|
|     Loop "Block2" (file k.cl line 10)
|     due to:
|     Data dependency on variable sum (file k.cl line 7)

serial region导致loop的执行成下图的状态：