Intel® FPGA SDK for OpenCL™: ベスト・プラクティス・ガイド

ID 683521
日付 12/08/2017
Public
ドキュメント目次

2.3.1. OpenCLのデザイン例のループ解析レポート

4つのループを含むOpenCLカーネルの例のループ解析レポート。

 1  // ND-Range kernel with unrolled loops
 2  __attribute((reqd_work_group_size(1024,1,1)))
 3  kernel void t (global int * out, int N) {
 4    int i = get_global_id(0);
 5    int j = 1;
 6    for (int k = 0; k < 4; k++) {
 7      #pragma unroll
 8      for (int n = 0; n < 4; n++) {
 9        j += out[k+n];
10      }
11    }
12    out[i] = j;
13
14    int m = 0;
15    #pragma unroll 1
16    for (int k = 0; k < N; k++) {
17      m += out[k/3];
18    }
19    #pragma unroll
20    for (int k = 0; k < 6; k++) {
21      m += out[k];
22    }
23    #pragma unroll 2
24    for (int k = 0; k < 6; k++) {
25      m += out[k];
26    }
27    out[2] = m;
28  }

このデザイン例のループ分析レポートは、コードで定義されているさまざまな種類のループのアンロール戦略を強調しています。

図 15. 4つのループを持つOpenCLデザイン例のループ解析レポート

Intel® FPGA SDK for OpenCL™オフライン・コンパイラーはソースコードに基づいて以下のループ・アンローリング戦略を実行します。

  • 最初のループ(ライン6)を完全に展開します。
  • #pragma unrollの仕様のため、最初のループ内(ライン8)で内側のループを完全に展開します。
  • #pragma unroll 1仕様のため、2番目の外部ループBlock2(ライン16)をアンロールしません。
  • #pragma unrollの仕様のため、3番目の外部ループ(ライン20)を完全に展開します。
  • #pragma unroll 2の仕様のため、4番目の外部ループBlock4(ライン24)を2回展開します。