Vitis™ ハードウェア アクセラレーション チュートリアル

xilinx.com の Vitis™ 開発環境を参照

モジュール 1 のコード、ファイル、および手順

このモジュール
1. ホスト/カーネルの実例の OpenCL API を確認
2. Vitis を GUI から実行、またはソフトウェアおよびハードウェア エミュレーションを実行する make を実行
3. Vitis アナライザーを実行してアプリケーション タイムラインを確認
4. Vitis HLS を実行してスループットのボトルネックを判別
5. HLS にテストベンチを挿入してすばやくイテレーションおよびカーネル コードの変更を検証

ホストとカーネルを使用したコード設定

各モジュールのソースコードは、ローカルの ./src ディレクトリの下にあります。

  • ホスト コードを確認すると、複数の API が使用されていることがわかります。データがカーネル間をどのように転送されるかに注目してください。実行モデルでは、次が実行されます。

    • ホスト プログラムが、Alveo データセンター アクセラレータ カード上の PCIe インターフェイスまたはエンベデッド プラットフォームの AXI バスを介して、カーネルで必要なデータをグローバル メモリに書き込みます。

    • ホスト プログラムが、入力パラメーターを使用してカーネルを設定します。

    • ホスト プログラムが FPGA のカーネル関数の実行をトリガーします。

    • カーネルが、必要に応じてグローバル メモリからのデータを読み出しながら、計算を実行します。

    • カーネルがグローバル メモリにデータを書き込み、ホストにタスクが終了したことを通知します。

    • ホスト プログラムがグローバル メモリからホスト メモリにデータを読み出し、必要に応じて処理を続けます。

  • カーネル コードを見てください。このコードは、Vitis ツールでコンパイルされ、ハードウェア記述 (Vivado ツールがザイリンクス デバイスにインプリメント) に変換されます。ホスト コードとカーネル コードは独立して開発およびコンパイルされるので、ヘッダー ファイルで extern “C” リンクを使用してカーネル関数宣言を囲むか、カーネル コードで関数全体を囲みます。

    extern "C" { void kernel_function(int *in, int *out, int size); }

Vitis での作成およびエミュレーション

このチュートリアルの各モジュールでは、GUI またはコマンド ライン (このチュートリアルの内容ではより効率的) を使用して Vitis を実行できます。

クリックして展開(Vitis GUI の手順)

Vitis を GUI で使用

ヒント: 次の画像を表示する場合、右クリックで別のタブやウィンドウで開くようにすると、内容を表示されたままにできます。

  1. ターミナルを開きます。

  2. Vitis を設定して起動します。

  3. [File] → [Import] をクリックします (画像)。

  4. Vitis プロジェクトでエクスポートした ZIP ファイルのデフォルトをそのまま使用し、[Next] をクリックします (画像)。

  5. 次のウィンドウで [Browse] をクリックし、./docs/module1_baseline/project に移動します (画像)。

  6. vitis_export_archive.ide.zip を選択し、[OK] をクリックします (画像)。

  7. 次のウィンドウで [Finish] をクリックします (画像)。

プラットフォームの設定

  1. 中央のウィンドウ ペインで、プラットフォーム リンクの直後にある 3 つのドット アイコン […] をクリックします (画像)。

  2. システムからアクセス可能なプラットフォームを指定します。

デザインのビルドとエミュレーション

[Assistant] ウィンドウに 3 つのメイン フローが表示されます。

  1. Emulation-SW: 論理レベルでデザインを検証します。

  2. Emulation-HW: カーネルをサイクル精度表記にコンパイルして、メトリクスをより正確に測定します。

  3. Hardware: フル コンパイルを開始し、ビットストリームを生成します。

ソフトウェア エミュレーションを実行します (数分かかります)。

ハードウェア エミュレーション実行します (約 10 ～ 20 分かかります)。

ハードウェア エミュレーションが完了したら、次の Vitis アナライザーのセクションに進みます。

または

クリックして展開(make の手順)

make の使用

  1. ターミナルを開きます。

  2. Vitis を設定します。

  3. ./build に移動します。

  4. make run TARGET=sw_emu を実行します (論理エミュレーションの場合)。

  5. make run TARGET=hw_emu を実行 (より詳細なエミュレーションの場合。10 分以上かかります)。

    • 現時点では hw ターゲットを実行すると 1 時間以上かかってしまうので、実行しないようにしてください。

Vitis アナライザーを使用したアプリケーションのエンド ツー エンドのタイムライン解析

Vitis アナライザーは、システム全体からカーネルの詳細まで、デザインのさまざまな側面を参照できるグラフィカル ツールです。

クリックして展開(Vitis Analyzer の手順)

  1. ターミナル、Vitis を設定します。

  2. vitis_analyzer & を実行します。

  3. [File] → [Open Summary] をクリックします。

  4. ./build を参照します。

  5. cholesky_kernel_hw_emu_xclbin_run_summary (青色の play アイコンが前に付いています)

  6. Vitis アナライザーの使用方法は、この 45 秒の looping gif を参照してください。

    次を確認してください。

    1. プロファイル サマリ

    2. ガイダンス レポート - 改善点を示します。

    3. アプリケーション タイムライン - 詳細は次を参照してください。

アプリケーション タイムラインの構造は次のとおりです。

  • ホスト

    • [OpenCL API Calls]: すべての OpenCL API 呼び出しがここでトレースされます。アクティビティ時間はホストの視点から測定されます。

    • [General]: clCreateProgramWithBinary、clCreateContext、および clCreateCommandQueue などの一般的な OpenCL API 呼び出しがここでトレースされます。

    • [Queue]: 特定のコマンド キューに関連する OpenCL API 呼び出しがここでトレースされます。これには clEnqueueMigrateMemObjects および clEnqueueNDRangeKernel などのコマンドが含まれます。ユーザー アプリケーションで複数のコマンド キューが作成された場合は、このセクションにすべてのキューとそのアクティビティが表示されます。

    • [Data Transfer]: このセクションでは、ホストからデバイス メモリまでの DMA 転送がトレースされます。OpenCL ランタイムにインプリメントされる DMA スレッドは複数あり、通常は同数の DMA チャネルがあります。DMA 転送は clEnqueueMigrateMemObjects などの OpenCLAPI を呼び出して、ユーザー アプリケーションにより開始されます。これらの DMA 要求がランタイムに転送され、スレッドの 1 つに割り当てられます。ホストからデバイスまでのデータ転送は [Write] の下、デバイスからホストまでのデータ転送は [Read] の下に表示されます。

    • [Kernel Enqueues]: ホスト プログラムによりエンキューされたカーネルが表示されます。ここに示されるカーネルを、デバイスのカーネル/CU と混同しないようにしてください。この「カーネル」は NDRangeKernels と、OpenCL コマンドの clEnqueueNDRangeKernels および clEnqueueTask で作成されるタスクのことです。これらはホストの視点から測定された時間に対してプロットされます。複数のカーネルが同時に実行されるようにスケジュールでき、実行がスケジュールされた時点からカーネル実行の終了までがトレースされます。重複するカーネル実行の数に応じて、複数のエントリが異なる行に表示されます。

  • [Device "name"]

    [Binary Container “name”]: バイナリ コンテナー名。

    • [Accelerator "name"]: FPGA 上の計算ユニット (アクセラレータ) の名前。

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

デバイス LUT およびフロップにインプリメントされるようになっている C++ カーネル (別名「ファブリック」) は、高位合成ツールの Vitis HLS で自動的にコンパイルされます。このチュートリアルでは、Vitis HLS を手動で実行して、基盤となる合成テクノロジとコレスキー カーネル アルゴリズムに関する追加情報を示します。

クリックして展開(Vitis HLS の手順)

  1. ターミナル、Vitis を設定します。

  2. ./build/cholesky_kernel_hw_emu/cholesky_kernel に移動します。

    • その階層には、もう 1 つ cholesky_kernel ディレクトリがあるはずです。

  3. vitis_hls -p cholesky_kernel & を実行して、Vitis 高位合成 GUI を起動します。

  4. Vitis HLS で高位合成レポートが表示されるようになりました。

  5. GUI で [Synthesis Summary Report] ウィンドウを展開します。

  6. [Performance & Resources] セクションでループと関数を展開します。

  7. このクリップの [II violation] を右クリックして、コード内の II 違反を見つけます。50s HLS looping GIF

注記: 元の Vitis HLS ウィンドウ レイアウトに戻すには、[Windows] メニューから [Reset Perspective] をクリックします。

開始間隔 (II)

この関数では、2 つのループに対する II 違反が 8 になっています。その 1 つは次のようになります。

// Loop only takes one element every 8 clock cycles!!!
// We expected one every cycle (II of 1)
for (int k = 0; k < j; k++)
{
    tmp += dataA[j][k] * dataA[j][k];
}

このバージョンのアルゴリズムでは、累算ありの double データ型を使用しているので、操作が実行されて終了するまでシリコンが 300MHz で 8 サイクルで動作する必要があります。サンプルは、1 つずつ 8 サイクルの間隔でしか計算できません。これが次のモジュールで取り組む最初のボトルネックです。

カーネル レイテンシ

それでは、レイテンシについて見てみましょう。

cholesky_kernel/solution/syn/report/cholesky_kernel_csynth.rpt

* Loop:
+--------------------+--------+---------+-------------+-----------+-----------+------------+----------+
|                    | Latency (cycles) |  Iteration  |  Initiation Interval  |    Trip    |          |
|       Loop Name    |  min   |   max   |   Latency   |  achieved |   target  |    Count   | Pipelined|
+--------------------+--------+---------+-------------+-----------+-----------+------------+----------+
|- VITIS_LOOP_32_..  |       ?|        ?|            3|          1|          1|           ?|    yes   |
|- Loop_first_col    |       ?|        ?|           34|          1|          1|           ?|    yes   |
|- Loop_col          |       ?|        ?|            ?|          -|          -|           ?|    no    |
| + Loop_diag        |      17|  2097161|           18|          8|          1| 1 ~ 262144 |    yes   |
| + Loop_row         |       ?|        ?| 61 ~ 2097205|          -|          -|           ?|    no    |
|  ++ Loop_vec_mul   |      17|  2097161|           18|          8|          1| 1 ~ 262144 |    yes   |
|- VITIS_LOOP_67_..  |       ?|        ?|            4|          1|          1|           ?|    yes   |
+--------------------+--------+---------+-------------+-----------+-----------+------------+----------+

次の点に注意してください。

  • VITIS が名前の先頭に付いたループ: Vitis HLS で自動的にラベル付けされたループで、ソース コードにはラベルは付いていません。表に含まれているそれ以外のループにはラベルが付いています。

  • クエスチョン マーク (?) は、関数へのスカラー入力に依存するために計算できないメトリクスを表します。この例では、行列サイズはコンフィギュレーション可能であり、レイテンシはそのサイズによって異なります。

  • 最後の Pipelined 列は、各サイクルで入力を処理するようにループが制約されているかどうかを示します。単純なループまたはほとんどの入れ子のループは、通常ツールによって自動的にパイプライン処理されます。

コレスキー関数への入力として、ユーザーが行列 N のサイズを渡します (この例の場合、64 でした)。

最初のループでは、開始間隔 II=1 で N 回の繰り返しが必要で、II=3 なので、終了するのに N x 3 かかります。Loop_first_col ループは N x 34 かかり、Loop_col ループは N 回 ((Loop_diag は N * 18) + (Loop_row は N * (N + 18)) 実行されます。最後のループは、最初のループと同様に N 回の繰り返しを必要とします。

次の場合、おおよその所要時間を見積もることができます。 N(18N+N(18N+residual1)+residual2) = 18N3 + (18+residual1)N2 + residual2.N

このため、基本的にアルゴリズムのレイテンシは、行列のサイズである N の 3 乗で計算できます。

カーネル用の C++ テストベンチの追加

このチュートリアルでは、カーネルをラップして Vitis HLS 環境でシミュレーションするあらかじめ作成された C++ の main プログラムを用意しています。

手順:

  1. ターミナルの docs ディレクトリから次を実行します。

    cp -r ./hls_tb ./module1_baseline/build/cholesky_kernel_hw_emu/cholesky_kernel
cp ./module1_baseline/src/cholesky_kernel.hpp ./module1_baseline/build/cholesky_kernel_hw_emu/cholesky_kernel/hls_tb

  2. Vitis HLS GUI が閉じている場合は、開き直します。

    cd ./module1_baseline/build/cholesky_kernel_hw_emu/cholesky_kernel
vitis_hls -p cholesky_kernel &

  3. GUI の [Explorer] ウィンドウの左ペインで [Source] の下の [Test Bench] を見つけます。右クリックして [Add file] を選択し、test_hls.cpp を選択します。この操作を ./hls_tb/tb_data の matrix_input.dat と golden_result.dat の 2 つのデータ ファイルに対して繰り返します。

  4. メインメニューで [Project] → [Run C simulation] をクリックします。これにより、「Csim」と呼ばれる純粋な論理シミュレーションが実行されます。HLS の合成するものは使用されません。

  5. [Project] → [Run C simulation] をクリックします。

  6. [Solution] → [Run C Synthesis] → [Active Solution] をクリックします。

  7. [Solution] → [Run C/RTL Cosimulation] をクリックします。ポップアップ ウィンドウで [OK] をクリックします。

Vitis HLS 協調シミュレーションは、クロック サイクルで実際レイテンシを示すサイクル精度の RTL シミュレーションを実行します。テストベンチでは、行列は 16x16 です。