パフォーマンスの最適化

ホスト最適化

このセクションでは、ホスト プログラムの最適化について説明し、OpenCL™ API を使用して、個別の計算ユニットの実行および FPGA とのデータ転送をスケジューリングします。このため、OpenCL コマンド キューを使用したタスクの並行実行について考慮する必要があります。このセクションでは、よくあるミスと、それらの見つけ方および解決方法を説明します。

カーネルをキューに追加するオーバーヘッドの削減

OpenCL API 実行モデルでは、データ並列およびタスク並列のプログラミング モデルがサポートされます。カーネルは通常 OpenCL ランタイムにより複数回エンキューされてから、デバイスで実行されるようにスケジュールされます。次のいずれかの方法でカーネルを開始するコマンドを送信する必要があります。

  • データ並列には clEnqueueNDRange API を使用
  • タスク並列には clEnqueueTask を使用

この送信プロセスはホスト プロセッサで実行されます。カーネルのコマンドおよび引数は、Alveo カードの場合は PCIe® バスを介して、アクセラレータに送信する必要があります。ザイリンクス ランタイム (XRT) ライブラリでは、アクセラレータへのコマンドおよび引数の送信のオーバーヘッドは、カーネルの引数の数によって 30 µs ~ 60 µs になります。このオーバーヘッドの影響は、カーネルを実行する必要のある回数を最小限にすることで削減ができます。

データ並列の場合、ザイリンクスでは、ホスト コードとカーネルのサイズに合わせて、グローバル ワーク サイズがローカル ワーク サイズの数倍になるようにサイズを選択することをお勧めします。理想的なのは、次のコード例のようにグローバル ワーク サイズとローカル ワーク サイズを同じにすることです。

size_t global = 1;
size_t local = 1;
clEnqueueNDRangeKernel(world.command_queue, kernel, 1, nullptr,
                       &global, &local, 2, write_events.data(),
                       &kernel_events[0]));
注記: タスク並列の場合、ザイリンクスでは clEnqueueTask の呼び出しを最小限にすることをお勧めします。理想的なのは、すべてのワークロードが clEnqueueTask の呼び出し 1 つで終了するようにすることです。

カーネル実行のオーバーヘッド削減の詳細は、カーネル実行 を参照してください。

データ移動の最適化

1: データ移動の最適化フロー

OpenCL API では、すべてのデータがまずホスト メモリからデバイスのグローバル メモリに転送され、グローバル メモリからカーネルに転送されて計算されます。この計算結果はカーネルからグローバル メモリに戻され、最後にグローバル メモリからホスト メモリに転送されます。カーネル最適化のストラテジを決定するには、どのようにすればデータを効率的に移動できるかを理解することが重要です。

注記: 計算を最適化する前に、アプリケーション内のデータ移動を最適化します。

データ移動の最適化では、計算が非効率であるとデータ移動がストールすることがあるので、データ転送コードを計算コードと分離することが重要です。ザイリンクスでは、この最適化段階でのみ、ホスト コードとカーネルをデータ転送コードを使用して変更することをお勧めします。目標は、PCIe の帯域幅および DDR 帯域幅を最大限に活用することにより、システム レベルのデータ スループットを最大にすることです。この目標を達成するには、通常ソフトウェア エミュレーション、ハードウェア エミュレーション、および FPGA での実行を何度か繰り返す必要があります。

データ転送とカーネル計算のオーバーラップ

データベース分析のようなアプリケーションでは、アクセラレーション デバイスで使用可能なメモリよりも大きなデータ セットが使用され、データ全体をブロック単位で転送して処理する必要があります。これらのアプリケーションで優れたパフォーマンスを達成するには、データ転送と計算をオーバーラップさせる手法が必要となります。

次は、GitHub の Vitis サンプル デザイン: 入門ホスト カテゴリにある overlap からの vadd カーネルのコード例です。

#define BUFFER_SIZE 256
#define DATA_SIZE 1024

//TRIPCOUNT indentifier
const unsigned int c_len = DATA_SIZE / BUFFER_SIZE;
const unsigned int c_size = BUFFER_SIZE;

extern "C" {
void vadd(int *c, int *a, int *b, const int elements) {
#pragma HLS INTERFACE m_axi port = c offset = slave bundle = gmem
#pragma HLS INTERFACE m_axi port = a offset = slave bundle = gmem
#pragma HLS INTERFACE m_axi port = b offset = slave bundle = gmem

#pragma HLS INTERFACE s_axilite port = c bundle = control
#pragma HLS INTERFACE s_axilite port = a bundle = control
#pragma HLS INTERFACE s_axilite port = b bundle = control
#pragma HLS INTERFACE s_axilite port = elements bundle = control
#pragma HLS INTERFACE s_axilite port = return bundle = control

    int arrayA[BUFFER_SIZE];
    int arrayB[BUFFER_SIZE];
    for (int i = 0; i < elements; i += BUFFER_SIZE) {
       #pragma HLS LOOP_TRIPCOUNT min=c_len max=c_len
        int size = BUFFER_SIZE;
        if (i + size > elements)
            size = elements - i;
    readA:
        for (int j = 0; j < size; j++) {
           #pragma HLS PIPELINE II=1
           #pragma HLS LOOP_TRIPCOUNT min=c_size max=c_size
            arrayA[j] = a[i + j];
        }

    readB:
        for (int j = 0; j < size; j++) {
           #pragma HLS PIPELINE II=1
           #pragma HLS LOOP_TRIPCOUNT min=c_size max=c_size
            arrayB[j] = b[i + j];
        }

    vadd_writeC:
        for (int j = 0; j < size; j++) {
           #pragma HLS PIPELINE II=1
           #pragma HLS LOOP_TRIPCOUNT min=c_size max=c_size
            c[i + j] = arrayA[j] + arrayB[j];
        }
    }
}
}

この例では、ホストで実行する必要のあるタスクは次の 4 つです。

  1. バッファー a の書き込み (Wa)
  2. バッファー b の書き込み (Wb)
  3. vadd カーネルを実行
  4. バッファー c の読み出し (Rc)

順不同コマンド キューを使用すると、次の図に示すように、データ転送とカーネル実行をオーバーラップできます。この例のホスト コードでは、カーネルが 1 セットのバッファーを処理している間に、ホストでもう 1 つのバッファーのセットを処理できるように、すべてのバッファーにダブル バッファリングが使用されます。

OpenCL event オブジェクトを使用すると、複雑な操作依存を簡単に設定して、ホスト スレッドとデバイス動作を同期できます。イベントは、操作のステータスを調べるための OpenCL オブジェクトです。イベント オブジェクトは、カーネル実行コマンド、メモリ オブジェクトに対する読み出し、書き込み、コピー コマンドにより作成されるか、clCreateUserEvent を使用して作成されたユーザー イベントです。これらのコマンドで返されるイベントをクエリすることにより、操作が完了したかどうかを確認できます。次の図の矢印は、最適なパフォーマンスを達成するために、イベント トリガーをどのように設定できるかを示しています。

2: イベント トリガーの設定

ホスト コードは、ループ内の 4 つのタスクをエンキューしてデータ セット全体を処理します。また、各タスクのデータ依存が満たされるように、異なるタスク間のイベント同期を設定します。ダブル バッファリングは、異なるメモリ オブジェクト値を clEnqueueMigrateMemObjects API に渡すことにより設定します。イベント同期は、各 API 呼び出しがほかのイベントを待ち、その API が終了してからそれ自身のイベントをトリガーするようにすると達成できます。

// THIS PAIR OF EVENTS WILL BE USED TO TRACK WHEN A KERNEL IS FINISHED WITH
// THE INPUT BUFFERS. ONCE THE KERNEL IS FINISHED PROCESSING THE DATA, A NEW
// SET OF ELEMENTS WILL BE WRITTEN INTO THE BUFFER.
vector<cl::Event> kernel_events(2);
vector<cl::Event> read_events(2);
cl::Buffer buffer_a[2], buffer_b[2], buffer_c[2];

for (size_t iteration_idx = 0; iteration_idx < num_iterations; iteration_idx++) {
    int flag = iteration_idx % 2;

    if (iteration_idx >= 2) {
        OCL_CHECK(err, err = read_events[flag].wait());
    }

    // Allocate Buffer in Global Memory
    // Buffers are allocated using CL_MEM_USE_HOST_PTR for efficient memory and
    // Device-to-host communication
    std::cout << "Creating Buffers..." << std::endl;
    OCL_CHECK(err,
                buffer_a[flag] =
                    cl::Buffer(context,
                                CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
                                bytes_per_iteration,
                                &A[iteration_idx * elements_per_iteration],
                                &err));
    OCL_CHECK(err,
                buffer_b[flag] =
                    cl::Buffer(context,
                                CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
                                bytes_per_iteration,
                                &B[iteration_idx * elements_per_iteration],
                                &err));
    OCL_CHECK(err,
                buffer_c[flag] = cl::Buffer(
                    context,
                    CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
                    bytes_per_iteration,
                    &device_result[iteration_idx * elements_per_iteration],
                    &err));

    vector<cl::Event> write_event(1);

    OCL_CHECK(err, err = krnl_vadd.setArg(0, buffer_c[flag]));
    OCL_CHECK(err, err = krnl_vadd.setArg(1, buffer_a[flag]));
    OCL_CHECK(err, err = krnl_vadd.setArg(2, buffer_b[flag]));
    OCL_CHECK(err, err = krnl_vadd.setArg(3, int(elements_per_iteration)));

    // Copy input data to device global memory
    std::cout << "Copying data (Host to Device)..." << std::endl;
    // Because we are passing the write_event, it returns an event object
    // that identifies this particular command and can be used to query
    // or queue a wait for this particular command to complete.
    OCL_CHECK(
        err,
        err = q.enqueueMigrateMemObjects({buffer_a[flag], buffer_b[flag]},
                                            0 /*0 means from host*/,
                                            NULL,
                                            &write_event[0]));
    set_callback(write_event[0], "ooo_queue");

    printf("Enqueueing NDRange kernel.\n");
    // This event needs to wait for the write buffer operations to complete
    // before executing. We are sending the write_events into its wait list to
    // ensure that the order of operations is correct.
    //Launch the Kernel
    std::vector<cl::Event> waitList;
    waitList.push_back(write_event[0]);
    OCL_CHECK(err,
                err = q.enqueueNDRangeKernel(
                    krnl_vadd, 0, 1, 1, &waitList, &kernel_events[flag]));
    set_callback(kernel_events[flag], "ooo_queue");

    // Copy Result from Device Global Memory to Host Local Memory
    std::cout << "Getting Results (Device to Host)..." << std::endl;
    std::vector<cl::Event> eventList;
    eventList.push_back(kernel_events[flag]);
    // This operation only needs to wait for the kernel call. This call will
    // potentially overlap the next kernel call as well as the next read
    // operations
    OCL_CHECK(err,
                err = q.enqueueMigrateMemObjects({buffer_c[flag]},
                                                CL_MIGRATE_MEM_OBJECT_HOST,
                                                &eventList,
                                                &read_events[flag]));
    set_callback(read_events[flag], "ooo_queue");

    OCL_CHECK(err, err = read_events[flag].wait());
}

次に示す [Application Timeline] ビューでは、計算ユニット vadd_1 が継続的に実行されており、データ転送時間は完全に隠されています。

3: データ転送時間が隠された [Application Timeline] ビュー

バッファー メモリの分割

メモリ バッファーの割り当ておよび割り当て解除により、DDR コントローラーでメモリが分割されることがあります。これにより、計算ユニットが論理的には並列実行できるはずなのに、最適なパフォーマンスが得られない可能性があります。

この問題は、異なる計算ユニットに対して複数の pthread が使用されており、スレッドでカーネルがエンキューされるたびに異なるサイズのデバイス バッファーが多数割り当てられ、解放される場合によく発生します。この場合、タイムライン トレースでカーネル実行間にギャップが表示され、プロセスがスリープ状態になっているように見えます。

ランタイムで割り当てられる各バッファーは、ハードウェアで連続している必要があります。大型メモリの場合、多数のバッファーの割り当ておよび割り当て解除が発生すると、その空間が空くのを待機する時間が長くなることがあります。これは、デバイス バッファーを割り当て、カーネルの異なるエンキュー間で再利用すると回避できます。

メモリの最適化の詳細は、バースト読み出しおよび書き込み を参照してください。

計算ユニットのスケジューリング

スケジューリング カーネルの動作は、全体的なシステム パフォーマンスに大きく影響します。これは、複数の計算ユニット (同じカーネルまたは別のカーネルのもの) をインプリメントする場合などは、さらに重要となってきます。このセクションでは、カーネルのスケジューリングに関連するさまざまなコマンド キューについて説明します。

複数の順序どおりのコマンド キュー

次の図に、2 つの順序どおりのコマンド キュー (CQ0 および CQ1) の例を示します。スケジューラは各キューからのコマンドを順序どおりに実行しますが、CQ0 および CQ1 からのコマンドはどの順序でも取り出すことができます。必要な場合は、CQ0 および CQ1 間の同期を管理する必要があります。

4: 2 つの順序どおりのコマンド キューの例

次は、concurrent_kernel_execution_c 例の host.cpp からのコードで、複数の順序どおりのコマンド キューを設定し、各キューにコマンドをエンキューしています。

    OCL_CHECK(err,
              cl::CommandQueue ordered_queue1(
                  context, device, CL_QUEUE_PROFILING_ENABLE, &err));
    OCL_CHECK(err,
              cl::CommandQueue ordered_queue2(
                  context, device, CL_QUEUE_PROFILING_ENABLE, &err));
...

    printf("[Ordered Queue 1]: Enqueueing scale kernel\n");
    OCL_CHECK(
        err,
        err = ordered_queue1.enqueueNDRangeKernel(
            kernel_mscale, offset, global, local, nullptr, &kernel_events[0]));

    set_callback(kernel_events[0], "scale");
...
    printf("[Ordered Queue 1]: Enqueueing addition kernel\n");
    OCL_CHECK(
        err,
        err = ordered_queue1.enqueueNDRangeKernel(
            kernel_madd, offset, global, local, nullptr, &kernel_events[1]));

    set_callback(kernel_events[1], "addition");
...
    printf("[Ordered Queue 2]: Enqueueing matrix multiplication kernel\n");
    OCL_CHECK(
        err,
        err = ordered_queue2.enqueueNDRangeKernel(
            kernel_mmult, offset, global, local, nullptr, &kernel_events[2]));
    set_callback(kernel_events[2], "matrix multiplication");

1 つの順不同コマンド キュー

次の図に、1 つの順不同コマンド キューの例を示します。スケジューラは、キューからのコマンドをどの順序でも実行できます。必要に応じて、ユーザーがイベントの依存性と同期を手動で定義する必要があります。

5: 1 つの順不同コマンド キュー

次は、concurrent_kernel_execution_c 例の host.cpp からのコードで、1 つの順不同コマンド キューを設定し、必要に応じてコマンドをエンキューしています。

    OCL_CHECK(
        err,
        cl::CommandQueue ooo_queue(context,
                                   device,
                                   CL_QUEUE_PROFILING_ENABLE |
                                       CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
                                   &err));

...
    printf("[OOO Queue]: Enqueueing scale kernel\n");
    OCL_CHECK(
        err,
        err = ooo_queue.enqueueNDRangeKernel(
            kernel_mscale, offset, global, local, nullptr, &ooo_events[0]));
    set_callback(ooo_events[0], "scale");
...
    // This is an out of order queue, events can be executed in any order. Since
    // this call depends on the results of the previous call we must pass the
    // event object from the previous call to this kernel's event wait list.
    printf("[OOO Queue]: Enqueueing addition kernel (Depends on scale)\n");

    kernel_wait_events.resize(0);
    kernel_wait_events.push_back(ooo_events[0]);

    OCL_CHECK(err,
              err = ooo_queue.enqueueNDRangeKernel(
                  kernel_madd,
                  offset,
                  global,
                  local,
                  &kernel_wait_events, // Event from previous call
                  &ooo_events[1]));
    set_callback(ooo_events[1], "addition");
...
    // This call does not depend on previous calls so we are passing nullptr
    // into the event wait list. The runtime should schedule this kernel in
    // parallel to the previous calls.
    printf("[OOO Queue]: Enqueueing matrix multiplication kernel\n");
    OCL_CHECK(err,
              err = ooo_queue.enqueueNDRangeKernel(
                  kernel_mmult,
                  offset,
                  global,
                  local,
                  nullptr, // Does not depend on previous call
                  &ooo_events[2]));
    set_callback(ooo_events[2], "matrix multiplication");

次の図に示す [Application Timeline] ビューは、複数の順序どおりのキューと 1 つの順不同キューの両方の方法を使用して、計算ユニット mmult_1 が計算ユニット mscale_1 および madd_1 と並列で実行されているところを示しています。

6: mult_1 が mscale_1 と madd_1 と並列実行されていることを示す [Application Timeline] ビュー

カーネル最適化

FPGA を使用する利点の 1 つは、特定のアルゴリズム用にカスタマイズしたデザインを作成できる柔軟性と機能です。これにより、アルゴリズムのスループットと消費電力をトレードオフするさまざまなインプリメンテーションを使用できます。次のガイドラインを使用すると、デザインの複雑性を制御して、必要なデザイン目標を達成しやすくなります。

カーネル計算の最適化

7: カーネル計算最適化のフロー

カーネル最適化は、カーネル インターフェイスにデータが到達したらすぐにすべてのデータを消費できるプロセッシング ロジックを作成することを目的として実行されます。主要なメトリクスは、開始間隔 (II) と呼ばれる、カーネルが新しい入力データを受信できるようになるまでのクロック サイクル数です。II の最適化は、これは通常、関数のパイプライン処理、ループ展開、配列分割、データフローなどの手法を使用してデータパスを一致させるようにプロセッシング コードを展開することによって達成されます。カーネル最適化の詳細は、カーネルのリンク を参照してください。

インターフェイス属性 (詳細なカーネル トレース)

詳細なカーネル トレースには、AXI トランザクションおよびそのプロパティが表示されます。AXI トランザクションは、グローバル メモリ側と AXI インターコネクトのカーネル側 ([Kernel "pass" 1:1:1]) に対して表示されます。次の図に、新しくアクセラレーションされたアルゴリズムの典型的なカーネル トレースを示します。

8: アクセラレーションされたアルゴリズムのカーネル トレース

パフォーマンスに関して注目すべきフィールドは、次のとおりです。

[Burst Length]
1 つのトランザクションで送信されるパッケージ数を示します。
[Burst Size]
1 つのパッケージの一部として転送されるバイト数を示します。

たとえば、[Burst Length] が 1 でパッケージごとに 4 バイトだとすると、ある程度の量のデータを転送するのに個別の AXI トランザクションが多く必要になります。

注記: Vitis コア開発キットでは、サイズが 4 バイト未満のバーストは、それより小さいデータが送信される場合でも作成されません。この場合、AXI バーストをイネーブルせずに連続のアイテムにアクセスすると、同じアドレスに対して複数の AXI 読み出しが見られることがあります。

そのため、バースト長が短く、バースト サイズが 512 ビットよりもかなり小さい場合、インターフェイス パフォーマンスを最適化できる可能性があります。

バースト データ転送の使用

データをバースト転送すると、メモリ アクセスのレイテンシが隠され、帯域幅の使用率およびメモリ コントローラーの効率が改善します。

注記: バースト転送は、連続したアドレス位置からの連続するデータ要求から推論されます。詳細は、バースト読み出しおよび書き込み を参照してください。

バースト転送が発生すると、詳細なカーネル トレースに表示されるバースト率とバースト長の値が大きくなります。

9: 詳細なカーネル トレースに表示されるバースト データ転送

上の図では、AXI インターコネクトの後のメモリ データ転送も異なる方法でインプリメントされているのがわかります (トランザクション時間が短縮)。これらのトランザクション上にカーソルを置くと、AXI インターコネクトが 16 x 4 バイトのトランザクションを 1 つの 1 x 64 バイトのパッケージ トランザクションにパックしたことがわかります。この方が、AXI4 帯域幅がより効率的に使用されます。次のセクションでは、この最適化手法について詳細に説明します。

バースト インターフェイスはコーディング スタイルとアクセス パターンによって大きく異なります。ただし、次のコード例に示すように、データ転送と計算を分離すると、バースト検出が容易になり、パフォーマンスが改善します。

void kernel(T in[1024], T out[1024]) {
    T tmpIn[1024];
    T tmpOu[1024];
    read(in, tmpIn);
    process(tmpIn, tmpOut);
    write(tmpOut, out);
}

つまり、read 関数が AXI 入力から内部変数 (tmpIn) に読み込みを実行し、計算は内部変数 tmpIn および tmpOut に対して演算を実行する process 関数でインプリメントされ、write 関数は生成された出力を取り込んで AXI 出力に書き込みます。

計算から読み出し関数と書き込み関数を分離すると、次のようになります。

  • 読み出し/書き込み関数の制御構造 (ループ) がシンプルになり、バースト検出がシンプルになります。
  • AXI インターフェイスから計算関数を分離すると、可能なカーネル最適化が単純になります。詳細は、カーネル最適化 を参照してください。
  • 内部変数はオンチップ メモリにマップされるので、AXI トランザクションよりも高速にアクセスできます。Vitis コア開発キットでサポートされるアクセラレーション プラットフォームには最大 10 MB のオンチップ メモリがあり、パイプ、ローカル メモリ、およびプライベート メモリとして使用できます。これらのリソースを効率的に使用することで、アプリケーションの効率およびパフォーマンスを大幅に向上できます。

全 AXI データ幅の使用

Vitis コンパイラでは、カーネル引数のデータ型に基づいてカーネルおよびメモリ コントローラー間のユーザー データ幅を設定できます。ザイリンクスでは、データ スループットを最大にするため、ユーザーがメモリ コントローラーの全データ幅にマップされるデータ型を選択することをお勧めします。サポートされるアクセラレーション カードすべてのメモリ コントローラーで 512 ビットのユーザー インターフェイスがサポートされており、これらは int16 または C/C++ 任意精度データ型 ap_int<512> などの OpenCL ベクター データ型にマップできます。

次の図では、バースト AXI トランザクション (バースト長 16) およびパッケージ サイズ 512 ビット (バースト サイズ 64 バイト) です。

10: バースト AXI トランザクション

この例は、AXI データ幅を最大にした良いインターフェイス設定と、実際のバースト トランザクションを示しています。

インターフェイスを宣言するのに使用される複雑な構造体またはクラスがあると、メモリ レイアウトやデータ パッケージの違いにより、ハードウェア インターフェイスが複雑になることがあります。これにより、複雑なシステムでデバッグするのが困難な問題が発生する可能性があります。

注記: カーネル引数には、32 ビット境界にパック可能なシンプルな構造体を使用することをお勧めします。構造体の使用に推奨される方法は、GitHub のザイリンクス オンボーディング例kernel_to_gmem カテゴリの Custom Data Type Example を参照してください。
OpenCL 属性を使用したデータ幅の設定

OpenCL API には、AXI データ幅の調整により自動化された方法をサポートする属性があります。前述のインターフェイス データ幅の変更は API でもサポートされますが、より大型の入力ベクターに対応するように、アルゴリズムのコードを C/C++ と同様に変更する必要があります。

コードを手動で変更しなくてもよいようにするには、次の OpenCL 属性を使用して、データ幅の拡張とアルゴリズムのベクター化を実行します。

次のような例があるとします。

__attribute__((reqd_work_group_size(64, 1, 1)))
__attribute__((vec_type_hint(int)))
__attribute__((xcl_zero_global_work_offset))
__kernel void vector_add(__global int* c, __global const int* a, __global const int* b) {
    size_t idx = get_global_id(0);
    c[idx] = a[idx] + b[idx];
}

この場合、ハード コード化されたインターフェイスは 32 ビット幅のデータパス (int *c, int* a, int *b) なので、直接インプリメントするとメモリ スループットが大幅に制限されますが、3 つの属性の値に基づいて、幅の自動拡張と変換が適用されます。

__attribute__((vec_type_hint(int)))
int が計算およびメモリ転送 (32 ビット) に主に使用されるデータ型であることを宣言します。これにより、AXI インターフェイスのターゲット帯域幅 (512 ビット) に基づいて、ベクター化/幅拡張の係数を計算できます。この例では係数は 16 = 512 ビット / 32 ビットで、理論上はベクター化が適用されると 16 個の値が処理されます。
__attribute__((reqd_work_group_size(X, Y, Z)))
ワーク アイテムの合計を定義します (XYZ は正の定数)。X*Y*Z はワーク アイテムの最大数であり、メモリ帯域幅が飽和する最大ベクター化係数を定義します。この例の場合、ワーク アイテムの合計は 64*1*1=64 です。

適用される実際のベクター化係数は、実際にコード記述されたデータ型または vec_type_hint で定義されるベクター化係数と、reqd_work_group_size で定義される可能な最大ベクター係数の最大公約数です。

可能な最大ベクター化係数を実際のベクター化係数で除算した商は、OpenCL 記述の残りのループ カウントとなります。このループはパイプライン処理されるので、パイプライン処理されたインプリメンテーションの利点を活かすため、複数のループ反復が残っていると有益な場合があります。特にベクター化された OpenCL コードに長いレイテンシがある場合に有益です。

__attribute__((xcl_zero_global_work_offset))
__attribute__((xcl_zero_global_work_offset)) は、ランタイムでグローバル オフセット パラメーターは使用されず、すべてのアクセスが揃っていることをコンパイラに指定します。これにより、ワーク グループのアライメントに関する有益な情報がコンパイラに伝わり、通常メモリ アクセスのアライメントに伝搬されます (ハードウェアがより少ない)。

これらの変換により、合成される実際のデザインが変わることに注意してください。部分的に展開されるループの場合、データが格納されるローカル配列の形状を変更する必要があります。これは通常問題なく動作しますが、まれに悪影響が出ることもあります。

次に例を示します。

  • 配列の分割で、分割係数を展開/ベクター化係数で除算できない。
    • このため、マルチプレクサーが多数必要となり、スケジューラで問題となります (メモリ使用量およびコンパイル時間が大幅に増加する可能性あり)。ザイリンクスでは、分割係数を 2 のべき乗にすることをお勧めします (ベクター係数は常に 2 のべき乗)。
  • ベクター化されるループに関係のないリソース制約があると、スケジューラで II が満たされないことを示すメッセージが表示される。
    • II は展開されたループで計算され、各反復のスループットが乗算されるので、パフォーマンスの低下につながるとは限りません。通常はそれでもパフォーマンスは向上します。
    • リソース制約がある可能性があり、これらの問題を解決するとパフォーマンスがさらに改善することを示すメッセージがスケジューラから表示されます。
    • ローカル配列は、通常ベクター化しない方法をでコードの後のセクションでアクセスされるので、自動的に再形成されません。

OpenCL パイプを使用したカーネル間通信のレイテンシの削減

OpenCL API 2.0 仕様には、パイプと呼ばれる新しいメモリ オブジェクトが導入されています。パイプには、FIFO として構成されたデータが格納されます。パイプ オブジェクトには、パイプから読み出してパイプに書き込むビルドイン関数を使用してのみアクセスできます。パイプ オブジェクトはホストからはアクセスできません。パイプを使用すると、データを外部メモリなしで FPGA 内の 1 つのカーネルから別のカーネルにストリーミングでき、全体的なシステム レイテンシを大幅に向上できます。詳細は、Khronos Group からの OpenCL C 仕様バージョン 2.0 の Pipe Functions を参照してください。

Vitis IDE では、パイプはすべてのカーネル関数の外部でスタティックに定義する必要があります。OpenCL 2.x clCreatePipe API を使用したダイナミック パイプ割り当てはサポートされていません。パイプの深さは、パイプ宣言内で OpenCL 属性 xcl_reqd_pipe_depth を使用して指定する必要があります。詳細は、xcl_reqd_pipe_depth を参照してください。

xcl_reqd_pipe_depth に指定されているように、有効な値は 16、32、64、128、256、512、1024、2048、4096、8192、16384、32768 です。

1 つのパイプは、異なるカーネル内に 1 つのプロデューサーおよびコンシューマーのみを持つことができます。
pipe int p0 __attribute__((xcl_reqd_pipe_depth(32)));

パイプには、ノンブロッキング モードの標準 OpenCL read_pipe() および write_pipe() ビルトイン関数、またはブロッキング モードのザイリンクス拡張 read_pipe_block() および write_pipe_block() 関数を使用してアクセス可能です。

パイプのステータスは、OpenCL get_pipe_num_packets() および get_pipe_max_packets() ビルトイン関数を使用してクエリできます。

次の関数シグネチャは現在サポートされているパイプ関数で、gentype はビルトイン OpenCL C スカラー整数または浮動小数点データ型を示します。
int read_pipe_block (pipe gentype p, gentype *ptr) 
int write_pipe_block (pipe gentype p, const gentype *ptr) 

次は GitHub の Xilinx Getting Started Examples からの dataflow/dataflow_pipes_ocl 例で、パイプを使用してブロッキング read_pipe_block() および write_pipe_block() 関数によりデータを 1 つの処理段階から次の処理段階に渡しています。

pipe int p0 __attribute__((xcl_reqd_pipe_depth(32)));
pipe int p1 __attribute__((xcl_reqd_pipe_depth(32)));
// Input Stage Kernel : Read Data from Global Memory and write into Pipe P0
kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void input_stage(__global int *input, int size)
{
    __attribute__((xcl_pipeline_loop)) 
    mem_rd: for (int i = 0 ; i < size ; i++)
    {
        //blocking Write command to pipe P0
        write_pipe_block(p0, &input[i]);
    }
}
// Adder Stage Kernel: Read Input data from Pipe P0 and write the result 
// into Pipe P1
kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void adder_stage(int inc, int size)
{
    __attribute__((xcl_pipeline_loop))
    execute: for(int i = 0 ; i < size ;  i++)
    {
        int input_data, output_data;
        //blocking read command to Pipe P0
        read_pipe_block(p0, &input_data);
        output_data = input_data + inc;
        //blocking write command to Pipe P1
        write_pipe_block(p1, &output_data);
    }
}
// Output Stage Kernel: Read result from Pipe P1 and write the result to Global
// Memory
kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void output_stage(__global int *output, int size)
{
    __attribute__((xcl_pipeline_loop))
    mem_wr: for (int i = 0 ; i < size ; i++)
    {
        //blocking read command to Pipe P1
        read_pipe_block(p1, &output[i]);
    }
}

[Device Traceline] ビューには、ハードウェア エミュレーション実行後の OpenCL パイプの詳細なアクティビティおよびストールが表示されます。この情報は、最適なアプリケーションのエリアおよびパフォーマンスを達成する正しい FIFO サイズを選択するために使用できます。

11: [Device Traceline] ビュー

計算の並列処理の最適化

デフォルトでは、C/C++ ではアルゴリズムは常に順に実行されるので、計算の並列処理は記述できません。OpenCL API では、ワーク グループに対して計算の並列処理を記述することはできますが、アルゴリズム記述内では並列処理を追加できません。FPGA のような完全にコンフィギャラブルな計算エンジンは柔軟性が高いので、計算の並列処理を試してみることができます。

データ並列処理のコード記述

アルゴリズムの FPGA へのインプリメンテーションで計算の並列処理を活用するには、まずソース コードの計算の並列処理が合成ツールで認識されるようにする必要があります。ループおよび関数は、ソース記述で計算の並列処理および計算ユニットを反映する主な候補ですが、ソース コードの構造によっては Vitis テクノロジで必要な変換を適用できないことがあるので、インプリメンテーションで計算の並列処理の利点が活かされているかどうかを検証することが重要です。

計算の並列処理には、ソース コードに反映されないものもあるので、その場合はソース コードに追加する必要があります。たとえば、カーネルは 1 つの入力値に演算を実行するよう記述されているのに、FPGA インプリメンテーションでは計算が効率的に複数の値に対して並列に実行されるようになることがあります。このような並列モデルについては、全 AXI データ幅の使用 を参照してください。512 ビットのインターフェイスは、int16 または C/C++ 任意精度データ型 ap_int<512> などの OpenCL ベクター データ型を使用して作成できます。

注記: これらのベクター型も、カーネル内のデータの並列処理をモデリングする方法として使用できます (int16 の場合は 16 個までのデータパスを並列処理可能)。ベクター型の使用に推奨される方法は、GitHub の Xilinx Getting Started Examplesvision カテゴリから Median Filter Example を参照してください。

ループの並列処理

ループは、アルゴリズム コードの繰り返しを示す基本的な C/C++/OpenCL API 手法です。次の例に、ループ構造のさまざまな側面を示します。

  for(int i = 0; i<255; i++) {
    out[i] = in[i]+in[i+1];
  }
  out[255] = in[255];

このコードは、最後の値を除き、配列のすべての値に対して実行され、2 つの連続する値を加算します。このループが記述どおりにインプリメントされると、ループの各反復に 2 サイクルかかるので、合計 510 サイクルかかります。この詳細は、HLS プロジェクトのスケジュール ビューアーで確認できます。

12: スケジュール ビューアーに表示されたインプリメント済みループ構造

Vivado 合成結果では、これは合計とレイテンシで示されます。

13: パフォーマンス見積もりの合成結果

ここで重要なのは、レイテンシ値と LUT の使用数です。たとえば、コンフィギュレーションによって、レイテンシが 511、LUT の使用数が 47 個になることがあります。これらの値は、インプリメンテーションでの選択によって大きく異なります。この例では、必要なエリアは小さいですが、レイテンシは長くなります。

ループの展開

ループを展開すると、モデルを完全に並列処理できるようになります。展開するループをマークしておくと、ツールで並列処理を最大限にできるようにインプリメンテションが作成されます。展開するループをマークするには、OpenCL ループに UNROLL 属性を指定します。

__attribute__((opencl_unroll_hint))

または、C/C++ ループで unroll プラグマを使用します。

#pragma HLS UNROLL

詳細は、ループ展開 を参照してください。

この例に適用すると、HLS プロジェクトのスケジュール ビューアーに次のように示されます。

14: スケジュール ビューアー

次の図に、パフォーマンスの見積もりを示します。

15: パフォーマンス見積もり

同じ計算を並列実行することにより、総レイテンシは大幅に改善されて 127 サイクルになり、計算ハードウェアは 4845 個の LUT に増加しています。

for ループを確認すると、各加算が前のループ反復とは完全に別なので、このアルゴリズムが 1 サイクルでインプリメントできないことがわかります。これは、out 変数にメモリ インターフェイスが使用されるからです。Vitis コア開発キットでは、配列に対してデュアル ポート メモリがデフォルトで使用されます。つまり、各サイクルでメモリに書き込むことができるのは、最大 2 つの値までだということです。このため、完全な並列インプリメンテーションを達成するには、次の例に示すように、out 変数をレジスタ内に保持する必要があります。

#pragma HLS array_partition variable= out complete dim= 0

詳細は、pragma HLS array_partition を参照してください。

この変換の結果は、スケジュール ビューアーで確認できます。

16: スケジュール ビューアーに表示された変換結果

この場合、見積もりは次のようになります。

17: 変換後のパフォーマンス見積もり

このコードは、組み合わせ関数としてインプリメントでき、何分の 1 かのサイクルで完了できます。

ループのパイプライン処理

ループのパイプライン処理 に説明されているように、ループをパイプライン処理すると、ループの反復を時間的にオーバーラップさせることができます。反復を同時に実行できるようにすると、リソースを反復間で共有でき (リソース使用量を削減)、展開されないループと比較して実行時間が短くなります。

パイプライン処理を C/C++ でイネーブルにするには、pragma HLS pipeline を使用します。

#pragma HLS PIPELINE

OpenCL API では、xcl_pipeline_loop 属性を使用します。

__attribute__((xcl_pipeline_loop))
注記: OpenCL API では、ループのパイプライン処理に xcl_pipeline_workitems を使用する方法もあります。ワーク アイテム ループは明示的に記述されないので、これらのループをパイプライン処理するにはこの属性が必要です。
__attribute__((xcl_pipeline_workitems))

この例の場合、HLS プロジェクトのスケジュール ビューアーの表示は次のようになります。

18: スケジュール ビューアーに表示されたパイプライン処理済みのループ

全体的な見積もりは次のようになります。

19: パフォーマンス見積もり

ループの各反復のレイテンシは 2 サイクルなので、オーバーラップする反復は 1 つだけです。これにより、総レイテンシは処理前の 1/2 の 257 サイクルになります。ループ展開よりも少ないリソースでレイテンシを削減できます。

ほとんどの場合、ループのパイプライン処理だけで全体的なパフォーマンスを改善できますが、パイプライン処理がどれだけ効率的かはループの構造によって異なります。一般的な制限事項は次のとおりです。

  • メモリ ポートまたはプロセス チャネルなどのようにリソースに限りがある場合、反復のオーバーラップ (開始間隔) が制限されます。
  • ループ運搬依存 (1 つの反復で計算された変数条件が次の反復に影響する) により、パイプラインの II が増加することがあります。

これらは高位合成中にレポートされ、スケジュール ビューアーで確認できます。最高のパフォーマンスを得るには、コードを修正してこれらの制限要素を取り除くか、依存性を取り除く (配列のメモリ インプリメンテーションを再構築、依存を完全になくすなど) ようにツールに命令する必要があります。

タスクの並列処理

タスクの並列処理を使用すると、データフロー並列処理の利点を活かすことができます。ループの並列処理とは異なり、タスクの並列処理ではタスク間で発生するバッファリングの利点を活かして、全実行ユニット (タスク) を並列実行できます。

次に例を示します。

void run (ap_uint<16> in[1024],
	  ap_uint<16> out[1024]
	  ) {
  ap_uint<16> tmp[128];
  for(int i = 0; i<8; i++) {
    processA(&(in[i*128]), tmp);
    processB(tmp, &(out[i*128]));
  }
}

このコードを実行すると、processA および processB 関数が順に 128 回実行されます。ループ内の processAprocessB のレイテンシが合わせて 278 だとすると、総レイテンシは次のように見積もられます。

20: パフォーマンス見積もり

余分なサイクルはループ設定が原因で、これはスケジュール ビューアーで確認できます。

C/C++ コードでは、タスクの並列処理は for ループに DATAFLOW プラグマを追加すると実行されます。

#pragma HLS DATAFLOW

OpenCL API コードでは、for ループ前に次の属性を追加します。

__attribute__ ((xcl_dataflow))

このトピックの詳細は、データフロー最適化HLS プラグマ、および OpenCL 属性 を参照してください。

HLS レポートの見積もりで示したように、タスク間にダブル (ピンポン) バッファーを使用すると、全体的なパフォーマンスを大幅に改善できます。

21: パフォーマンス見積もり

この場合、異なる反復で異なるタスクが並行実行されるので、デザインの全体的なレイテンシがほぼ 1/2 になります。各関数の処理に 139 サイクルかかり、128 反復が完全にオーバーラップしているので、総レイテンシは次のようになります。

(1x only processA + 127x both processes + 1x only processB) * 139 cycles = 17931 cycles 

タスクの並列処理は、インプリメンテーションでパフォーマンスを改善できる効果的な手法ですが、DATAFLOW プラグマを任意のコードに適用した場合にどれくらい効果的かは大きく異なります。DATAFLOW プラグマの最終的なインプリメンテーションを理解するには、個々のタスクの実行パターンを確認することが必要です。Vitis コア開発キットでは、同時実行を示す詳細なカーネル トレースが提供されています。

22: 詳細なカーネル トレース

上の図に示すように、この詳細なカーネル トレースでは、データフロー フロー ループの始点が表示されます。プロセス A はループの最初に即座に開始し、プロセス B はプロセス A の終了を待ってから最初の反復を開始します。プロセス B がループの最初の反復を完了している間に、プロセス A は 2 回目の反復の演算を開始します。

この情報のより抽象的な表示は、ホストとデバイスのアクティビティの アプリケーション タイムライン に示されます。

計算ユニットの最適化

データ幅

パフォーマンスに関しては、インプリメンテーションに必要なデータ幅が重要な要素の 1 つです。ツールはアルゴリズム全体にポート幅を伝搬します。アルゴリズム記述から開始した場合は特に、C/C++/OpenCL API コードで、デザインのポートにも、整数型などの大型データ型のみが使用されることがあります。ただし、アルゴリズムが完全にコンフィギャラブルなインプリメンテーションにマップされていくと、10 ビットまたは 12 ビットなどのより小型のデータ型で十分なこともあります。このため、最適化中に HLS 合成レポートで基本的な演算のサイズを確認することをお勧めします。

通常は、Vitis コア開発キットでアルゴリズムを FPGA にマップする際、C/C++/OpenCL API 構造を理解て動作依存を抽出するため、多くの処理が必要になります。このマップを実行するため、Vitis 開発キットによりソース コードが演算ユニットに分割され、FPGA にマップされます。これらの演算ユニット (ops) の数およびサイズには、さまざまな要素が影響します。

次の図では、基本的な演算とそのビット幅がレポートされています。

23: 演算の使用量の見積もり

アルゴリズム記述でよく使用される典型的なビット幅 (16、32、64 ビット) を探して、C/C++/OpenCL API ソースからの関連する演算にこれほどのビット幅が本当に必要なのかを検証します。演算が小さいほど計算時間も短くなるので、これによりアルゴリズムのインプリメンテーションが大幅に改善する可能性があります。

固定小数点の演算

アプリケーションによっては、ほかのハードウェア アーキテクチャ用に最適化されているというだけの理由で、浮動小数点計算が使用されているものがあります。深層学習のようなアプリケーションに固定小数点演算を使用すると、精度を同程度に保ちながら、消費電力とエリアを大幅に節約できます。

注記: ザイリンクスでは、浮動小数点演算を使用することを確定する前に、アプリケーションに固定小数点演算を使用することを検討してみることをお勧めします。

マクロ演算

より大型の計算エレメントを考慮すると有益な場合もあります。ツールでは、ソース コードが残りのソース コードとは別に実行され、アルゴリズムが周囲の演算を考慮せずに FPGA にマップされます。この場合、Vitis テクノロジで演算の境界が維持され、特定のコードに対して実質的にマクロ演算が作成されます。これには、次の原則が使用されます。

  • マップ プロセスに対する演算局所性
  • 経験則のための複雑性の削減

これにより、結果が大きく異なるものになることがあります。C/C++ では、マクロ演算は #pragma HLS inline off を使用すると作成されます。 OpenCL API では、関数を定義する際に次の属性を指定しないことにより、同様のマクロ演算を生成できます。

__attribute__((always_inline))

詳細は、pragma HLS inline を参照してください。

最適化済みライブラリの使用

OpenCL 仕様には、多数の数学ビルトイン関数が含まれます。native_ 接頭辞が付いた数学ビルトイン関数はすべて 1 つまたは複数のネイティブ デバイス命令にマップされ、通常は対応する関数 (native_ 接頭語なし) よりも優れたパフォーマンスになります。これらの関数の精度と入力範囲 (場合による) は、インプリメンテーションで定義されます。Vitis テクノロジでは、これらの native_ ビルトイン関数に対して、ザイリンクス FPGA 用にエリアおよびパフォーマンスを最適済みの Vivado HLS ツールの Math ライブラリに含まれる同等の関数が使用されます。

注記: ザイリンクスでは、精度がアプリケーション要件を満たす場合は、native_ ビルトイン関数または HLS ツールの Math ライブラリを使用することをお勧めします。

メモリ アーキテクチャの最適化

メモリ アーキテクチャはインプリメンテーションの重要な側面です。帯域幅のアクセスには制限があり、次の図に示すように全体的なパフォーマンスに大きく影響することがあります。


void run (ap_uint<16> in[256][4],
          ap_uint<16> out[256]
         ) {
  ...
  ap_uint<16> inMem[256][4];
  ap_uint<16> outMem[256];

  ... Preprocess input to local memory
  
  for( int j=0; j<256; j++) {
    #pragma HLS PIPELINE OFF
    ap_uint<16> sum = 0;
    for( int i = 0; i<4; i++) {

      sum += inMem[j][i];
    }
    outMem[j] = sum;
  } 

  ... Postprocess write local memory to output
}

このコードでは、2 次元入力配列の内部次元に関連する 4 つの値が追加されます。これ以上変更をしないでインプリメントすると、次のような見積もりになります。

24: パフォーマンス見積もり

全体的なレイテンシが 4608 (Loop 2) なのは、18 サイクル (内部ループ 16 サイクル + 合計のリセット + 書き出される出力) が 256 回反復されているからです。これは、HLS プロジェクトのスケジュール ビューアーで確認できます。見積もりは、内部ループを展開すると大幅に改善されます。

25: パフォーマンス見積もり

ただし、このように改善されるのは、主にプロセスがデュアル ポート メモリの両方のポートを使用しているからです。これは、プロジェクトのスケジュール ビューアーから確認できます。

26: スケジュール ビューアー

メモリからのすべての値にアクセスして合計を計算するために、2 つの読み出しがサイクルごとに実行されています。これにより、メモリへのアクセスが完全にブロックされてしまうので、望ましくない結果になることがあります。結果をさらに改善するには、2 番目の次元を使用してメモリを 4 つの小型メモリに分割します。

#pragma HLS ARRAY_PARTITION variable=inMem complete dim=2

詳細は、pragma HLS array_partition を参照してください。

これにより、4 つの配列読み出しになり、すべてが 1 つのポートを使用して異なるメモリで実行されます。

27: 4 つの配列の実行結果

Loop 2 に合計 256 x 4 サイクル = 1024 サイクルが使用されます。

28: パフォーマンス見積もり

または、メモリを 4 ワードの 1 つのメモリに再形成します。これには、次のプラグマを使用します。

#pragma HLS array_reshape variable=inMem complete dim=2

詳細は、pragma HLS array_reshape を参照してください。

これにより、配列分割と同じレイテンシになりますが、この場合は 1 つのポートを使用した 1 つのメモリになります。

29: レイテンシ結果

どちらのソリューションでも全体的なレイテンシおよび使用量は同じようになりますが、配列を再形成した方がインターフェイスがきれいで、配線の密集は少なくなります。

注記: これで配列最適化は終了です。実際のデザインでは、ループを並列処理するとレイテンシがさらに改善できることがあります (ループの並列処理 を参照)。
void run (ap_uint<16> in[256][4],
	  ap_uint<16> out[256]
	  ) {
  ...

  ap_uint<16> inMem[256][4];
  ap_uint<16> outMem[256];
  #pragma HLS array_reshape variable=inMem complete dim=2
  
  ... Preprocess input to local memory
  
  for( int j=0; j<256; j++) {
    #pragma HLS PIPELINE OFF
    ap_uint<16> sum = 0;
    for( int i = 0; i<4; i++) {
      #pragma HLS UNROLL
      sum += inMem[j][i];
    }
    outMem[j] = sum;
  } 

  ... Postprocess write local memory to output

}

カーネル SLR および DDR メモリの割り当て

デザインの周波数およびリソース要件を満たすには、カーネル計算ユニット (CU) インスタンスおよび DDR メモリ リソースのフロアプランが重要となります。フロアプランでは、CU (カーネル インスタンス) を明示的に SLR に割り当てたり、CU を DDR メモリ リソースにマップしたりします。フロアプランする際、CU のリソース使用率と DDR メモリの帯域幅の要件を考慮してください。

最大のザイリンクス FPGA は複数のスタックド シリコン ダイで構成されています。各スタックは SLR (Super Logic Region) と呼ばれ、DDR インターフェイスなど、決まった量のリソースおよびメモリが含まれます。カスタム ロジックに使用可能なデバイス SLR リソースについては、Vitis 2019.2 ソフトウェア プラットフォーム リリース ノート を参照してください。または、platforminfo ユーティリティ で説明する platforminfo ユーティリティを使用して表示することもできます。

実際のカーネル リソース使用率を使用して CU を複数の SLR に分配すると、特定の SLR での密集を削減できます。システム見積もりレポートでは、デザイン サイクルの早期に、カーネルで使用される多くのリソース (LUT、フリップフロップ、BRAM など) を確認できます。このレポートは、コマンド ラインまたは GUI を使用して、ハードウェア エミュレーションおよびシステム コンパイル中に生成できます。詳細は、システム見積もりレポート を参照してください。

この情報と使用可能な SLR リソースの情報を使用して、1 つの SLR が過剰に使用されないように、CU を SLR に割り当てます。SLR の密集が少ないほど、ツールでデザインを FPGA リソースに適切にマップしやすくなり、パフォーマンス ターゲットを満たすことができます。メモリ リソースと CU のマップについては カーネル ポートのグローバル メモリへのマップ および 計算ユニットの SLR への割り当て を参照してください。

注記: 計算ユニットは使用可能な DDR メモリ リソースのいずれにでも接続できますが、SLR に割り当てる際は、カーネルの帯域幅要件を考慮する必要もあります。

CU を SLR に割り当てたら、CU マスター AXI ポートを DDR メモリ リソースにマップします。ザイリンクスでは、CU と同じ SLR にある DDR メモリ リソースに接続することをお勧めします。そのようにすると、数が決まっている SLR をまたぐ接続リソースの競合を削減できます。また、SLR 間の接続には SLL (Super Long Line) 配線リソースが使用されるので、標準の SLR 内の配線よりも遅延が大きくなります。

SLR 領域をまたいで別の SLR にある DDR リソースに接続することが必要なことはありますが、connectivity.sp および connectivity.slr 指示子の両方が明示的に指定されている場合は、ツールで自動的にクロッシング ロジックが追加され、SLL 遅延の影響を最小限に抑えて、タイミング クロージャが達成されるようになっています。

複数のメモリ バンクにアクセスするカーネルのガイドライン

DDR メモリ リソースは、プラットフォームの SLR (Super Logic Region) をまたいで分配されます。SLR をまたぐ接続の数は制限されるので、カーネルを DDR メモリ リソースとの接続数が最も多い SLR に配置するのが一般的です。これにより、SLR をまたぐ接続の競合が削減し、SLR をまたぐためにロジック リソースが消費されるのを回避できます。

30: 同じ SLR 内のカーネルおよびメモリ
注記: 左の図では、1 つの AXI インターフェイスが 1 つのメモリ バンクに接続されています。右の図では、複数の AXI インターフェイスが同じメモリ バンクに接続されています。

上の図に示すように、カーネルに 1 つの AXI インターフェイスがあり、1 つのメモリ バンクにのみマップされる場合、platforminfo ユーティリティ で説明される platforminfo ユーティリティにより、カーネルのメモリ バンクに接続されている SLR がリストされるので、これがカーネルが最適に配置される SLR です。この場合、追加の入力なしでもデザイン ツールによりカーネルが自動的にその SLR に配置される可能性はありますが、次の場合は明示的に SLR を割り当てる必要があります。

  • デザインに同じメモリ バンクにアクセスするカーネルが多数含まれる場合。
  • カーネルにメモリ バンクの SLR に含まれない特殊なロジック リソースが必要な場合。

カーネルに複数の AXI インターフェイスがあり、すべてのインターフェイスが同じメモリ バンクにアクセスする場合、1 つの AXI インターフェイスを含むカーネルと同様に処理でき、カーネルを AXI インターフェイスがマップされているメモリ バンクと同じ SLR に配置する必要があります。

31: 隣接する SLR のメモリ バンク
注記: 左の図では、カーネルが SLR0 にあり、SLR をまたぐ接続が 1 つ必要です。右の図では、カーネルがメモリ バンクにアクセスするために、SLR をまたぐ接続が 2 つ必要です。

カーネルに複数の AXI インターフェイスがあり、異なる SLR にある複数のメモリ バンクに接続されている場合は、カーネルがアクセスするメモリ バンクの大部分が含まれる SLR にカーネルを配置することをお勧めします。これにより、このカーネルで必要な SLR をまたぐ接続数が最小限に抑えられ、ユーザー デザイン内のほかのカーネルでメモリ バンクに接続するために使用可能な SLR をまたぐリソースが増えます。

カーネルが別の SLR にあるメモリ バンクに接続される場合は、カーネル SLR および DDR メモリの割り当て のように SLR の割り当てを明示的に指定します。

32: 2 つ離れた SLR のメモリ バンク
注記: 左の図では、マップされているすべてのメモリ バンクにアクセスするために SLR をまたぐ接続が 2 つ必要です。右の図では、マップされているすべてのメモリ バンクにアクセスするために SLR をまたぐ接続が 3 つ必要です。

プラットフォームに含まれる SLR が 3 つ以上になると、上の図に示すように、最も多くマップされるメモリ バンクのすぐ隣ではない SLR のメモリ バンクにカーネルがマップされることもあります。このような場合、離れたメモリ バンクにアクセスするために複数の SLR 境界をまたぐ必要があるので、SLR をまたぐリソースの使用量が増加します。このようなリソース使用量の増加を避けるには、カーネルを中央の SLR に配置して、隣接する SLR にまたぐリソースのみが使用されるようにします。

Vivado HLS を使用したカーネルの最適化

OpenCL または C/C++ を使用したカーネルの最適化は、Vitis コア開発キットから実行できます。このセクションで説明されているような主な最適化制約 (関数およびループのパイプライン処理、関数およびループ間で同時処理を増加するためのデータフローの適用、ループの展開など) は、Vivado HLS ツールで実行されます。

Vitis コア開発キットは、自動的に HLS ツールを呼び出しますが、GUI の解析機能を使用するには、Vitis テクノロジ内から HLS ツールを直接起動する必要があります。Vivado HLS でのカーネルのコンパイル に説明されているように、スタンドアロン モードで HLS ツールを使用すると、最適化を次のように実行できます。

  • エミュレーションを実行する必要はないので、カーネルの最適化のみに集中できます。
  • 複数のソリューションを作成し、結果を比較し、ソリューション スペースを調べて、最適なデザインを見つけることができます。
  • インタラクティブな [Analysis] パースペクティブを使用してデザイン パフォーマンスを解析できます。
重要: Vitis コア開発キットに戻すのはカーネル ソース コードのみです。最適化スペースを調べたら、すべての最適化がカーネル ソース コードに OpenCL 属性または C/C++ プラグマとして適用されるようにします。

HLS ツールをスタンドアロン モードで開くには、[Assistant] ビューでハードウェア関数オブジェクトを右クリックし、Open HLS Project をクリックします (次の図を参照)。

33: HLS プロジェクトを開く

トポロジ最適化

このセクションでは、トポロジ最適化について、複数の計算ユニットの大まかなレイアウトとインプリメンテーションに関連する属性と、パフォーマンスへの影響について説明します。

複数の計算ユニット

ターゲット デバイスで使用可能なリソースによって、同じカーネルまたは異なるカーネルの複数の計算ユニットを作成して並列で実行することで、システムの処理時間とスループットを改善できます。詳細は、複数のカーネル インスタンスの作成 を参照してください。

複数 DDR バンクの使用

Vitis テクノロジでサポートされるアクセラレーション カードには、最大 80 GB/s の生 DDR 帯域幅の 1、2、または 4 つの DDR バンクが含まれます。FPGA と DDR の間で大量のデータを移動するカーネルの場合、ザイリンクスでは Vitis コンパイラおよびランタイム ライブラリで複数の DDR バンクを使用するように指示することをお勧めします。

ホスト アプリケーションは、DDR バンクだけでなく、カーネルに直接データを転送する PLRAM にアクセスできます。この機能をイネーブルにするには、設定ファイルに connnectivity.sp オプションを含め、v++ --config でその設定ファイルを指定します。この最適化の使用法の詳細は カーネル ポートのグローバル メモリへのマップ、グローバル メモリ バンクへのデータ転送に関する詳細は メモリマップドのインターフェイス を参照してください。

複数の DDR バンクの利点を活かすには、ホスト コードで CL メモリ バッファーを異なるバンクに割り当て、xclbin ファイルを v++ コマンド ラインでのバンク割り当てと同じになるように設定する必要があります。

次の図に、GitHub の Vitis Examples からの Global Memory Two Banks (C) 例のブロック図を示します。この例では、カーネルの入力ポインター インターフェイスを DDR バンク 0 に、出力ポインター インターフェイスを DDR バンク 1 に接続しています。

34: グローバル メモリの 2 つのバンクの例

ホスト コードでの DDR バンクの割り当て

ホスト コードでのバンク割り当ては、ザイリンクス ベンダー拡張でサポートされています。次に、必要なヘッダー ファイルからの抜粋と、入力バッファーを DDR バンク 0 に、出力バッファーをバンク 1 にそれぞれ割り当てるコード例を示します。

#include <CL/cl_ext.h>
…
int main(int argc, char** argv) 
{
…
    cl_mem_ext_ptr_t inExt, outExt;  // Declaring two extensions for both buffers
    inExt.flags  = 0|XCL_MEM_TOPOLOGY; // Specify Bank0 Memory for input memory
    outExt.flags = 1|XCL_MEM_TOPOLOGY; // Specify Bank1 Memory for output Memory
    inExt.obj = 0   ; outExt.obj = 0; // Setting Obj and Param to Zero
    inExt.param = 0 ; outExt.param = 0;

    int err;
    //Allocate Buffer in Bank0 of Global Memory for Input Image using Xilinx Extension
    cl_mem buffer_inImage = clCreateBuffer(world.context, CL_MEM_READ_ONLY | CL_MEM_EXT_PTR_XILINX,
            image_size_bytes, &inExt, &err);
    if (err != CL_SUCCESS){
        std::cout << "Error: Failed to allocate device Memory" << std::endl;
        return EXIT_FAILURE;
    }
    //Allocate Buffer in Bank1 of Global Memory for Input Image using Xilinx Extension
    cl_mem buffer_outImage = clCreateBuffer(world.context, CL_MEM_WRITE_ONLY | CL_MEM_EXT_PTR_XILINX,
            image_size_bytes, &outExt, NULL);
    if (err != CL_SUCCESS){
        std::cout << "Error: Failed to allocate device Memory" << std::endl;
        return EXIT_FAILURE;
    }
…
}

cl_mem_ext_ptr_tstruct で、次のように定義されます。

typedef struct{
    unsigned flags;
    void *obj;
    void *param;
  } cl_mem_ext_ptr_t;
  • flags の有効な値は、次のとおりです。
    • XCL_MEM_DDR_BANK0
    • XCL_MEM_DDR_BANK1
    • XCL_MEM_DDR_BANK2
    • XCL_MEM_DDR_BANK3
    • <id> | XCL_MEM_TOPOLOGY
      注記: <id> は xxx.xclbin ファイルの次に生成される xxx.xclbin.info ファイルの [Memory Configuration] セクションから判断されます。xxx.xclbin.info ファイルには、グローバル メモリ (DDR、PLRAM など) が <id> を示すインデックスと共にリストされます。
  • obj: CL_MEM_USE_HOST_PTR フラグが clCreateBuffer API に渡された場合に CL メモリ バッファーに割り当てられるホスト メモリに関連付けられているポインターです。それ以外の場合は NULL に設定されます。
  • param: 今後の使用のために予約されています。常に 0 または NULL に設定します。

カーネル コードのためのグローバル メモリの割り当て

複数の AXI インターフェイスの作成

OpenCL カーネル、C/C++ カーネル、および RTL カーネルでは、AXI インターフェイスへの関数パラメーターの割り当て方法はそれぞれ異なります。

  • OpenCL カーネルでは、カーネル引数の各グローバル ポインターに AXI4 インターフェイスを 1 つ生成するのに --max_memory_ports オプションが必要です。AXI4 インターフェイス名は、引数リストのグローバル ポインターの順序に基づいて付けられます。

    次のコード例は、GitHub の Vitis Getting Started Exampleskernel_to_gmem カテゴリにある gmem_2banks_ocl 例からのものです。

    __kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
    void apply_watermark(__global const TYPE * __restrict input, 
    __global TYPE * __restrict output, int width, int height) {
     ...
    }

    この例では、1 つ目のグローバル ポインター inputAXI4M_AXI_GMEM0 を、2 つ目のグローバル ポインター outputM_AXI_GMEM1 を割り当てています。

  • C/C++ カーネルでは、異なるグローバル ポインターの HLS INTERFACE プラグマに異なる bundle 名を指定することで、複数の AXI4 インターフェイスが生成されます。詳細は、カーネル インターフェイスおよびメモリ バンク を参照してください。

    次の gmem_2banks_c からのコード例では、input ポインターをバンドル gmem0 に、output ポインターをバンドル gmem1 に割り当てています。バンドル名には任意の有効な C 文字列を使用でき、生成される AXI4 インターフェイス名は M_AXI_<bundle_name> になります。この例の場合、入力ポインターの AXI4 インターフェイス名は M_AXI_gmem0、出力ポインター名は M_AXI_gmem1 になります。詳細は、pragma HLS interface を参照してください。
    #pragma HLS INTERFACE m_axi port=input  offset=slave bundle=gmem0
    #pragma HLS INTERFACE m_axi port=output offset=slave bundle=gmem1
    
  • RTL カーネルでは、RTL Kernel ウィザードでのインポート プロセスでポート名が生成されます。RTL Kernel ウィザードで付けられるデフォルト名は m00_axi および m01_axi です。変更しない場合は、これらの名前を設定ファイルの connectivity.sp オプションで DDR バンクを割り当てるときに使用する必要があります。詳細は、カーネル ポートのグローバル メモリへのマップ を参照してください。
DDR バンクへの AXI インターフェイスの割り当て
重要: ザイリンクスでは、DDR インターフェイスを複数使用する場合は、各カーネル/CU に DDR メモリ バンクを指定し、カーネルを配置する SLR を配置することを要件としています。詳細は、カーネル ポートのグローバル メモリへのマップ および 計算ユニットの SLR への割り当て を参照してください。

次に、connectivity.sp オプションを指定する設定ファイルの例と、入力ポインター (M_AXI_GMEM0) を DDR バンク 0 に、出力ポインター (M_AXI_GMEM1) を DDR バンク 1 に接続する v++ のコマンド ライン例を示します。

config_sp.txt ファイル:

[connectivity] 
sp=apply_watermark_1.m_axi_gmem0:DDR[0] 
sp=apply_watermark_1.m_axi_gmem1:DDR[1]

v++ コマンド ライン:

v++ apply_watermark --config config_sp.txt

[Device Hardware Transaction] ビューを使用すると、実際の DDR バンクの通信を確認して DDR の使用を解析できます。

35: [Device Hardware Transaction] ビューの DDR バンクのトランザクション
PLRAM への AXI インターフェイスの割り当て

一部のプラットフォームでは、PLRAM がサポートされます。その場合、DDR バンクへの AXI インターフェイスの割り当て で説明されているのと同じ --connectivity.sp オプションを使用しますが、名前は PLRAM[id] を使用します。特定のプラットフォームでサポートされる有効な名前は、xclbin と共に生成される xclibin.info ファイルの [Memory Configuration] セクションに表示されます。

カーネルの SLR 領域への割り当て

グローバル メモリ バンクにポートを割り当てるには、割り当てられた DDR、HBM、またはブロック RAM に接続するため、カーネルを FPGA に物理的に配線する必要あります。大型の FPGA では、複数の SLR (Super Logic Region) を含むスタックド シリコン デバイスが使用されています。Vitis コア開発キットでは、計算ユニットはターゲット プラットフォームと同じ SLR に配置されます。特にカーネルが別の SLR 領域にある特定のメモリ バンクに接続されている場合など、これが望ましくない場合もあります。その場合は、カーネル インスタンス (CU) を手動でグローバル メモリと同じ SLR に割り当てます。詳細は、カーネル ポートのグローバル メモリへのマップ を参照してください。

CU インスタンスを SLR に割り当てるには、計算ユニットの SLR への割り当てで説明されている connectivity.slr オプションを使用します。

ヒント: DDR および SLR 領域の数などのプラットフォーム属性について理解するには、platforminfo ユーティリティ に説明されている platforminfo コマンドを使用してターゲット プラットフォームの詳細を表示できます。