アプリケーションのプロファイリング

Vitis™ コア開発キットでは、コンパイル中にシステムおよびカーネルのリソースとパフォーマンスに関するさまざまなレポートが生成されます。これらのレポートは、アプリケーションのパフォーマンスのベースラインを確立し、ボトルネックを特定して、ハードウェアでアクセラレーション可能なターゲット関数を決定 (デバイス アクセラレーション アプリケーションの設計手法 を参照) するのに役立ちます。ザイリンクス ランタイム (XRT) では、エミュレーションおよびシステム モードのコンフィギュレーションの両方で、アプリケーションの実行中にプロファイリング データも収集されます。次のようなデータがレポートされます。

  • ホストおよびデバイスのタイムライン イベント
  • OpenCL™ API の呼び出しシーケンス
  • カーネルの実行シーケンス
  • AXI トランザクションを含む FPGA トレース データ
  • カーネルの開始信号および停止信号

これらのレポートとプロファイリング データは、アプリケーションのパフォーマンスのボトルネックを特定し、デザインを最適化してパフォーマンスを向上するために使用できます。アプリケーションの最適化には、アプリケーション ホスト コードとハードウェアでアクセラレーションされるカーネルの両方の最適化が必要です。ホスト コードはデータ転送とカーネル実行がスムーズに実行されるように最適化する必要があり、カーネルではパフォーマンスとリソース使用量が適切なものになるよう最適化する必要があります。

Vitis でアルゴリズムを最適化する際には、システムのリソース使用量とパフォーマンス、カーネル最適化、ホスト最適化、およびデータ転送最適化の 4 つのエリアがあります。次の Vitis レポートおよびグラフィカル ツールは、これらのエリアをプロファイリングおよび最適化するのに役立ちます。

レポートは、コマンド ライン (アプリケーションの実行 を参照) または Vitis 統合設計環境 (IDE) のどちらから実行しても、アクティブ ビルドを実行したときに自動的に生成されます。3 つのビルド ターゲットすべてに対して、異なるレポート セットが生成され、該当するレポート ディレクトリに保存されます。これらのレポートのディレクトリに関する詳細は、ディレクトリ構造を参照してください。

レポートは、ウェブ ブラウザー、スプレッドシート ビューアー、または Vitis GUI で表示できます。Vitis IDE でこれらのレポートを開くには、[Assistant] ビューでレポートをダブルクリックします。

この後のセクションで、さまざまなレポートおよびグラフィカル表示ツールについて説明し、デザインのプロファイリングに使用する方法を示します。

機能とパフォーマンスのベースライン

最適化を開始する前に、まずアプリケーションのパフォーマンスを理解しておくことが重要です。これには、アプリケーションの機能とパフォーマンスのベースラインを確立します。

1: 機能とパフォーマンスのベースライン フロー

ボトルネックの検出

まず、ターゲット プラットフォームで実行しているアプリケーションのボトルネックを特定します。最も効果的なのは、valgrindcallgrind、および GNU gprof などのプロファイリング ツールを使用してアプリケーションを実行する方法です。これらのツールで生成されるプロファイリング データには、すべての関数呼び出しの回数と実行時間を示すコールグラフが表示されます。

ソフトウェアおよびハードウェア エミュレーションの実行

アプリケーションの実行 に説明されているように、アクセラレーション アプリケーションのソフトウェアおよびハードウェア エミュレーションを実行して機能が正しいかどうかを確認し、ホスト コードおよびカーネルのプロファイリング データを生成します。カーネル コンパイル レポート、プロファイル サマリ、タイムライン トレース、およびデバイス ハードウェア トランザクションを解析して、タイミング間隔、レイテンシ、DSP やブロック RAM などのリソース使用量のベースライン パフォーマンス見積もりを理解します。

アプリケーションのビルドおよび実行

ベースライン作成の最後に、アプリケーションをビルドして FPGA アクセラレーション カード (Alveo™ データセンター アクセラレータ カードなど) で実行します。システム コンパイルからのレポートおよびアプリケーション実行からのプロファイリング データを解析し、ハードウェアでの実際のパフォーマンスおよびリソース使用量を確認します。

ヒント: ベースライン作成中のレポートはすべて保存して、最適化時に参照して結果を比較できるようにしておきます。

設計ガイダンス

Vitis コア開発キットには、デザインで検出された問題に対して、ソフトウェア アプリケーション開発者が即座に実行可能なガイダンスを提供する包括的な設計ガイダンス ツールがあります。問題の原因は、ソース コードに関係している場合や、実行されていないツールの最適化である場合があります。また、ルールはさまざまな基準デザイン セットからの経験に基づく汎用ルールであり、特定のデザインには適用しないこともあります。そのため、ガイダンス ルールを理解して、特定のアルゴリズムおよび要件に基づいて適切な操作を実行するようにしてください。

ガイダンスは、Vivado HLS、Vitis プロファイラー、および v++ コンパイラで起動された Vivado Design Suite から生成されます。生成された設計ガイダンスには、エラー、アドバイザリ、警告、およびクリティカル警告の重要度があり、ソフトウェア エミュレーション、ハードウェア エミュレーション、システム ビルド中に表示されます。プロファイルの設計ガイダンスは、プロファイリング結果を解釈し、パフォーマンスを向上するのに役立ちます。

ガイダンスには、ハイパーリンク、例、および資料へのリンクが含まれています。これにより、問題をすばやく特定し、Vitis テクノロジの使用に関する情報にアクセスできるので、生産性が向上します。

設計ガイダンスは、Vitis IDE でデザインをビルドまたは実行すると自動的に生成されます。

ガイダンス レポートを開く方法は、Vitis アナライザーの使用 に説明されています。ガイダンス レポートにアクセスするには、コンパイル サマリ、リンク サマリ、または実行サマリを開き、ガイダンス レポートを開きます。ガイダンス レポートを開いたら、ガイダンスの上にカーソルを置くと、推奨されるソリューションが表示されます。

次の図に、Vitis アナライザーで表示されるガイダンス レポートの例を示します。リンクをクリックすると、拡張ビューに実行可能なガイダンスが表示されます。

2: 設計ガイダンスの例

v++ コマンドの各実行 (コンパイルおよびリンクを含む) に対して、1 つの HTML ガイダンス レポートが生成されます。レポート ファイルは、--report_dir の特定の出力名の下に保存されます。次に例を示します。

  • v++ コンパイル: v++_compile_<output>_guidance.html
  • v++ リンク: v++_link_<output>_guidance.html

ガイダンス レポートの表示

カーネルがコンパイルされたとき、および FPGA バイナリがリンクされたときに、v++ コマンドによりガイダンス レポートが自動的に生成されます。これらのレポートを表示するには、Vitis アナライザーからアプリケーション プロジェクトの <output_filename>.compile.summary または <output_filename>.link.summary を開きます。<output_filename>v++ コマンドの出力です。

たとえば、Vitis アナライザーを起動してレポートを開くには、次のコマンドを使用します。

vitis_analyzer <output_filename>.link.summary

Vitis アナライザーが開くと、リンク サマリ レポート、コンパイル サマリ、およびコンパイルおよびリンク プロセスで生成されたさまざまなレポートが表示されます。左側のペインで Build をクリックすると、コンパイルとリンクの両方の段階でガイダンス レポートが生成されます。詳細は、Vitis アナライザーの使用 を参照してください。

ガイダンス データの解釈

[Guidance] ビューには、入力項目が行ごとに表示されます。各行には、ガイダンス ルール名、しきい値、実際の値、そのルールの簡単な説明が表示されます。最後のフィールドには、ルール違反を理解して回避するために役立つ参照資料へのリンクが含まれます。

GUI の [Guidance] ビューでは、[Name] 列にイダンス ルールと ID がカテゴリ別にまとめられ、重要度を示すシンボルと共に表示されます。これらは、HTML レポートに個別にリストされます。また、HTML レポートにはヒントは表示されませんが、[Full Name] 列が含まれます。

次に、HTML ガイダンス レポートに含まれるすべてのフィールドとその説明を示します。

Id
各ガイダンス ルールには ID 値が割り当てられています。この ID は、ガイダンス レポートから特定のメッセージを見つけるために使用します。
Name
ガイダンス ルールを識別するためのニーモニック名が表示されます。これらの名前は、特定のガイダンス ルールを記憶しやすくするために付けられています。
Severity
ガイダンス ルールの重要度を示します。
Full Name
[Name] 列のニーモニック名よりも詳細な名前が示されます。
Categories
ほとんどのメッセージはカテゴリに分類されています。これにより、GUI で [Guidance] ビューの共通ツリー ノードの下にメッセージが論理的なカテゴリ別に表示されます。
Threshold
ルールが満たされているかどうかの判断基準となるしきい値が表示されます。しきい値は、良いデザインおよびコーディング プラクティスに従っている多数のアプリケーションに基づいて決定されています。
Actual
デザインでの実際の値が表示されます。この値が基準値と比較され、ルールが満たされたかどうかが判断されます。
Details
現在のルールの詳細を説明するメッセージが表示されます。
Resolution
ルールを満たすために変更可能なモデル ソース コードまたはツール変更によく使用される方法へのリンクが含まれます。リンクをクリックすると、特定の問題に適用可能なヒントやコードを含むポップアップ ウィンドウまたは資料が開きます。

システム見積もりレポート

プロセスで実行時間が最も長いのは、ハードウェア システムのビルドとザイリンクス デバイス上での FPGA バイナリの実行です。ビルド時間は、ターゲット デバイスと FPGA ファブリックにインスタンシエートする計算ユニット数にもよります。そのため、システム ハードウェア用にビルドせずにアプリケーションのパフォーマンスを見積もることができれば便利です。

システム見積もりレポートには、FPGA リソース使用量およびハードウェアでアクセラレーションされたカーネルが動作可能な見積もり周波数が示されます。このレポートは、ハードウェア エミュレーションおよびシステム ハードウェア ビルドで自動的に生成されます。レポートには、リソース使用量および見積もり周波数を含むユーザー カーネルの全体的な情報が含まれます。このレポートをデサイン最適化のガイドとして使用できます。

システム見積もりレポートは、次のオプションを使用して強制的に生成することもできます。

v++ .. --report_level estimate

次の図に、サンプル レポートを示します。

3: システム見積もり

システム見積もりリレポートを開く

システム見積もりレポートは、Vitis アナライザー ツールで開くことができ、Vitis コンパイラでアプリケーションをビルドしたとき、および XRT ライブラリでアプリケーションを実行したときに表示されるようになっています。Vitis アナライザーを起動してレポートを開くには、次のコマンドを使用できます。

vitis_analyzer <output_filename>.link.summary

<output_filename>v++ コマンドの出力です。これにより Vitis アナライザー ツールでアプリケーション プロジェクトのリンク サマリが開くので、[System Estimate] レポートを選択します。詳細は、Vitis アナライザーの使用 を参照してください。

ヒント: システム見積もりレポートはテキスト ファイルなので、テキスト エディターまたはターゲット プラットフォームでも表示できます。

システム見積もりレポートの解釈

v++ で生成されるシステム見積もりレポートには、アプリケーションのバイナリ コンテナーごとおよびデザインの計算ユニットごとの情報が示されます。レポートの構成は次のとおりです。

  • ターゲット デバイス情報
  • アプリケーションのカーネルごとのサマリ
  • バイナリ コンテナーごとの詳細な情報

次に、見積もりレポートに含まれる情報の例を示します。

-------------------------------------------------------------------------------
Design Name:             mmult.hw_emu.xilinx_u200_xdma_201830_2
Target Device:           xilinx:u200:xdma:201830.2
Target Clock:            300.000000MHz
Total number of kernels: 1
-------------------------------------------------------------------------------

Kernel Summary
Kernel Name  Type  Target              OpenCL Library                          Compute Units
-----------  ----  ------------------  --------------------------------------  -------------
mmult        c     fpga0:OCL_REGION_0  mmult.hw_emu.xilinx_u200_xdma_201830_2  1


-------------------------------------------------------------------------------
OpenCL Binary:     mmult.hw_emu.xilinx_u200_xdma_201830_2
Kernels mapped to: clc_region

Timing Information (MHz)
Compute Unit  Kernel Name  Module Name  Target Frequency  Estimated Frequency
------------  -----------  -----------  ----------------  -------------------
mmult_1       mmult        mmult        300.300293        411.015198

Latency Information (clock cycles)
Compute Unit  Kernel Name  Module Name  Start Interval  Best Case  Avg Case  Worst Case  
------------  -----------  -----------  --------------  ---------  --------  ----------  
mmult_1       mmult        mmult        826 ~ 829       825        827       828         

Area Information
Compute Unit  Kernel Name  Module Name  FF     LUT    DSP   BRAM  URAM
------------  -----------  -----------  -----  -----  ----  ----  ----
mmult_1       mmult        mmult        81378  35257  1036  2     0
-------------------------------------------------------------------------------

デザインおよびターゲット デバイスのサマリ

すべてのデザイン見積もりレポートの最初に、ターゲット デバイスに関するアプリケーションのサマリと情報が表示されます。デバイス情報は、レポートのその後のセクションに含まれます。

-------------------------------------------------------------------------------
Design Name:             mmult.hw_emu.xilinx_u200_xdma_201830_2
Target Device:           xilinx:u200:xdma:201830.2
Target Clock:            300.000000MHz
Total number of kernels: 1
-------------------------------------------------------------------------------

デザイン サマリには、次の情報が含まれます。

Target Device
Vitis コンパイラでビルドされた FPGA バイナリを実行するターゲット プラットフォームに含まれるザイリンクス デバイスの名前。
Target Clock
FPGA ファブリックにマップされた計算ユニットの動作周波数。

カーネル サマリ

このセクションには、アプリケーション プロジェクトに対して定義されたカーネルすべてがリストされます。次に、カーネル サマリの例を示します。

Kernel Summary
Kernel Name  Type  Target              OpenCL Library                          Compute Units
-----------  ----  ------------------  --------------------------------------  -------------
mmult        c     fpga0:OCL_REGION_0  mmult.hw_emu.xilinx_u200_xdma_201830_2  1

カーネル名に加え、実行ターゲットおよび入力ソースのタイプも表示されます。OpenCL™、C、C/C++ ソース ファイルではコンパイルおよび最適化手法が違うので、カーネル ソース ファイルのタイプが指定されます。

[Kernel Summary] セクションは、レポートの最後に表示されるサマリ情報で、この後に各計算ユニットのバイナリ コンテナーに関する詳細な情報が含まれます。

タイミング情報

各バイナリ コンテナーの詳細セクションには、各計算ユニット (CU) の実行ターゲットと、タイミング情報が示されます。通常、FPGA バイナリの周波数見積もりがターゲット周波数よりも高い場合は、CU はデバイスで実行できます。周波数見積もりがターゲット周波数よりも低い場合は、FPGA ファブリックで正しく実行されるように、CU のカーネル コードをさらに最適化がする必要があります。次に、この情報の例を示します。

OpenCL Binary:     mmult.hw_emu.xilinx_u200_xdma_201830_2
Kernels mapped to: clc_region

Timing Information (MHz)
Compute Unit  Kernel Name  Module Name  Target Frequency  Estimated Frequency
------------  -----------  -----------  ----------------  -------------------
mmult_1       mmult        mmult        300.300293        411.015198

重要なのは、ターゲット周波数と周波数見積もりの差を理解することです。CU は FPGA ファブリック内に隔離して配置されているわけではありません。CU は、アプリケーションのクラスをサポートするためにデバイス開発者が定義したその他のコンポーネントを含む有効な FPGA デザインの一部として配置されます。

CU カスタム ロジックはカーネルごとに生成されるので、見積もり周波数がターゲットよりも高いということは、CU が高い見積もり周波数で実行できるということです。そのため、FPGA バイナリのインプリメンテーションでは、CU がターゲット周波数でタイミングを満たす必要があります。

レイテンシ情報

レイテンシ情報には、バイナリ コンテナーの各 CU の実行プロファイルが示されます。このデータを解析する際は、すべての値が CU 境界からカスタム ロジックを介してされていることに注意してください。グローバル メモリへのデータ転送に関連するインシステム レイテンシは、これらの値の一部としてはレポートされません。また、FPGA ファブリックでターゲットとなる CU のみのこのレイテンシがレポートされます。次に、レイテンシ レポートの例を示します。

Latency Information (clock cycles)
Compute Unit  Kernel Name  Module Name  Start Interval  Best Case  Avg Case  Worst Case  
------------  -----------  -----------  --------------  ---------  --------  ----------  
mmult_1       mmult        mmult        826 ~ 829       825        827       828         

レイテンシ レポートは、次のフィールドに分けられます。

  • 開始間隔
  • ベスト ケース レイテンシ
  • 平均ケース レイテン
  • ワースト ケース レイテンシ

開始間隔は、特定のカーネルにおける CU の実行間の時間を定義します。

ベスト、平均、ワースト ケース レイテンシは、CU がそのカーネルの 1 つの ND Range データ タイルの結果を生成するのにかかる時間を示します。カーネルにデータが依存する計算ループがない場合、レイテンシは同じになります。ループ実行にデータ依存があると、データ特定のレイテンシ変動が発生し、これがレイテンシ レポートに含まれます。

開始間隔およびレイテンシは、次の条件のカーネルに対しては「undef」 (未定義) とレポートされます。
  • OpenCL カーネルに明示的な reqd_work_group_size(x,y,z) がない
  • カーネルに可変境界のループがある
注記: レイテンシ情報は、ループの変換とそのモデルの並列処理の解析に基づいた見積もりを反映します。パイプライン処理およびデータフローなどの高度な変換があると、実際のスループットが大きく変わります。このため、レイテンシは実行間の相対的な比較にのみ使用してください。

エリア情報

FPGA は空白の計算キャンバスと考えることもできますが、各 FPGA で使用可能な基本的な構築ブロック数には限りがあります。これらの基本的な構築ブロック (FF、LUT、DSP、ブロック RAM) は、Vitis コンパイラでデザイン内の計算ユニットごとにカスタム ロジックを生成するために使用されます。1 つの CU のカスタム ロジックにインプリメントする必要のある基本的なリソース数により、FPGA ファブリックに同時に読み込むことができる CU の数が決まります。次に、1 つの CU に対してレポートされるエリア情報を示します。

Area Information
Compute Unit  Kernel Name  Module Name  FF     LUT    DSP   BRAM  URAM
------------  -----------  -----------  -----  -----  ----  ----  ----
mmult_1       mmult        mmult        81378  35257  1036  2     0
-------------------------------------------------------------------------------

HLS レポート

HLS レポートは、ハードウェア エミュレーションおよびシステム ビルドで生成され、ユーザー カーネルの高位合成 (HLS) プロセスに関する詳細を提供します。このプロセスでは、カーネル ロジックを FPGA にインプリメントするため、C/C++ および OpenCL カーネルをハードウェア記述言語に変換します。カスタム ハードウェア ロジックの FPGA リソース使用量の見積もり、動作周波数、レイテンシ、およびインターフェイス信号が示されます。これらの詳細は、カーネルの最適化に役立つ多数の情報を提供します。

コマンド ラインから実行すると、このレポートは次のディレクトリに保存されます。

_x/<kernel_name>.<target>.<platform>/<kernel_name>/<kernel_name>/solution/syn/report

HLS レポートを表示するには、Vitis アナライザーの使用 に説明されているように、Vitis アナライザーからコンパイル サマリまたはリンク サマリを開きます。次に、HLS レポートの例を示します。

4: HLS レポート

HLS レポートの生成と表示

重要: Vivado HLS で生成されるレポートを含む中間ファイルを保存するには、--save-temps オプションを指定しする必要があります。HLS レポートおよび HLS ガイダンスは、C および OpenCL カーネルのハードウェア エミュレーションおよびシステム ビルドでのみ生成されます。ソフトウェア エミュレーションおよび RTL カーネルには生成されません。

HLS レポートを表示するには、Vitis アナライザーからアプリケーション プロジェクトの <output_filename>.compile.summary または <output_filename>.link.summary を開きます。<output_filename>v++ コマンドの出力です。

Vitis アナライザーを起動してレポートを開くには、次のコマンドを使用できます。

vitis_analyzer <output_filename>.compile.summary

Vitis アナライザーが開くと、コンパイルのサマリと、コンパイル プロセスで生成されたレポートが表示されます。HLS レポートを表示するには、左側のペインで Build をクリックます。詳細は、Vitis アナライザーの使用 を参照してください。

HLS レポートの解釈

HLS レポートの左側のペインには、モジュール階層が表示されます。この階層には、HLS 実行の一部として生成された各モジュールが示されます。これらのモジュールのいずれかを選択すると、[Synthesis Report] ウィンドウの右側にそのモジュールの合成の詳細が表示されます。合成レポートは、複数のセクションから構成されています。

  • 一般情報
  • パフォーマンス見積もり (タイミングおよびレイテンシ)
  • 使用率見積もり
  • インターフェイス情報

この情報が階層ブロックの一部である場合は、その階層に含まれるブロックの情報すべてがまとめられます。どのインスタンスがデザイン全体に影響しているのかがわかっている場合は、レポート内から階層をナビゲートすることもできます。

注意: サイクルおよびレイテンシの絶対数に関しては、これらの値は HLS 合成時の見積もりに基づいているので、パイプラインおよびデータフローなどの高度な変換を使用する場合は特に、最終的な結果を正確に反映していない可能性があります。レポートにクエスチョン マーク (?) が表示される場合、原因は可変境界ループにある可能性があるので、このようなループにはトリップカウントを設定し、このレポートに相対的な見積もりが表示されるようにすることをお勧めします。

プロファイル サマリ レポート

正しく設定されていれば、Vitis ランタイム ライブラリでホスト アプリケーションおよびカーネルのプロファイリング データが収集されます。アプリケーションの実行が終了すると、プロファイル サマリが .csv ファイルとしてホスト コードがコンパイルされたディレクトリに保存されます。

プロファイル サマリには、全体的なアプリケーション パフォーマンスに関する注釈付きの詳細が表示されます。アプリケーションの実行中に生成されたすべてのデータが複数のカテゴリに分類されます。プロファイル サマリでは、カーネル実行とデータ転送の統計を確認できます。

ヒント: プロファイル サマリ レポートは、すべてのビルド コンフィギュレーションで生成できます。ただし、ソフトウェア エミュレーション ビルドでは、カーネル実行効率およびデータ転送効率の下にデータ転送の詳細は含まれません。この情報は、ハードウェア エミュレーションおよびシステム ビルド コンフギュレーションでのみ生成されます。

次に、プロファイル サマリ レポートの例を示します。

5: [Profile Summary]

レポートには複数のタブがあります。次の表に、各タブの説明を示します。

[Top Operations]
カーネルおよびグローバル メモリ。最上位演算のサマリを示します。FPGA とデバイス メモリ間の最上位データ転送のプロファイル データを表示します。
[Kernels & Compute Units]
すべてのカーネルおよび計算ユニットのプロファイル データを表示します。
[Data Transfers]
ホストおよびグローバル メモリ。ホストとデバイス メモリ間の PCIe リンクを介したすべての読み出しおよび書き込み転送のプロファイル データを表示します。カーネルとグローバル メモリ間のデータ転送がイネーブルの場合、その情報も表示します。
[OpenCL APIs]
ホスト アプリケーションで実行されるすべての OpenCL C ホスト API 関数のプロファイル データを表示します。
[Kernel Internals]
このタブは、xrt.ini ファイル に示すように、xrt.ini[Emulation] セクションで launch_waveform をイネーブルにすると、ハードウェア エミュレーション中に表示されます。生成される波形データ (.wdb) は profile_kernels.csv および timeline_kernels.csv ファイルでレポートされ、この情報を示す [Kernel Internals] タブが表示されます。レポートされる情報は C/C++ および OpenCL カーネルに該当するものです。RTL カーネルに対してはレポートされません。

プロファイル サマリの詳細は、プロファイル サマリの解釈 を参照してください。

プロファイル サマリ レポートの生成と表示

プロファイル サマリに必要なデータをキャプチャするには、アプリケーションを実行する前に次の手順を実行する必要があります。

  1. FPGA バイナリ (xclbin) ファイルは、デフォルトでプロファイリング データをキャプチャするよう設定されます。ただし、リンク プロセス中に v++ --profile_kernel オプションを使用すると、キャプチャしたプロファイリング データが詳細に表示できるようになります。--profile_kernel オプションの詳細は、Vitis コンパイラ コマンド を参照してください。
  2. ランタイムでは、xrt.ini ファイル に説明されているように、xrt.ini ファイルにプロファイリング データをキャプチャするためのキーワードを含める必要があります。
    [Debug]
    profile = true
  3. [Kernel Internals] データのプロファイリングをイネーブルにするには、xrt.ini[Emulation] セクションに launch_waveform タグも追加する必要があります。
    [Emulation]
    launch_waveform = batch

FPGA バイナリおよび xrt.ini ファイルでプロファイリングをイネーブルにすると、アプリケーションを実行したときに、ランタイムで profile_summary.csv レポート ファイルが生成されるほか、[Kernel Internals] をオンにしたときは profile_kernels.csv および timeline_kernels.csv ファイルも作成されます。

CSV レポートは、スプレッドシート ツールまたはユーティリティ、あるいは Vitis アナライザー ツールで開くことができ、アプリケーションをビルドしたときに Vitis コンパイラから、アプリケーションを実行したときに XRT ライブラリから表示されるようになっています。Vitis アナライザーを起動してレポートを開くには、次のコマンドを使用できます。

vitis_analyzer profile_summary.csv

関連情報

プロファイル サマリの解釈

プロファイル サマリには、OpenCL アプリケーションに役立つさまざまな統計が含まれ、アプリケーションの機能的なボトルネックの概要が表示されます。プロファイル サマリには、4 つのセクションから構成され、次の情報が表示されます。

  1. Top Operations
    • Top Data Transfer: Kernels and Global Memory: FPGA とデバイス メモリ間の最上位データ転送のプロファイル データを表示します。
      • Device: デバイス名
      • Compute Unit: 計算ユニット名
      • Number of Transfers: デバイスで監視される書き込みおよび読み出し AXI トランザクションの合計
      • Average Bytes per Transfer: (読み出しバイト数合計 + 書き込みバイト数合計) / (読み出し AXI トランザクション数合計 + 書き込み AXI トランザクション数合計)
      • Transfer Efficiency (%): (転送ごとの平均バイト数) / min(4K, (メモリ ビット幅/8 * 256))

        AXI4 仕様では、最大バースト長が 256、最大バースト サイズが 4K バイトに制限されています。

      • Total Data Transfer (MB): (読み出しバイト数合計 + 書き込みバイト数合計) / 1.0E6
      • Total Write (MB): (書き込みバイト数合計) / 1.0E6
      • Total Read (MB): (読み出しバイト数合計) / 1.0E6
      • Transfer Rate (MB/s): (データ転送数合計) / (計算ユニットの合計時間)
    • Top Kernel Execution
      • Kernel Instance Address: カーネル インスタンスのホスト アドレス (16 進数)
      • Kernel: カーネル名
      • Context ID: ホストのコンテキスト ID
      • Command Queue ID: ホストのコマンド キュー ID
      • Device: カーネルが実行されたデバイスの名前 (フォーマット:<device>-<ID>)
      • Start Time (ms): 実行の開始時間 (ms)
      • Duration (ms): 実行期間 (ms)
      • Global Work Size: カーネルの NDRange
      • Local Work Size: カーネルのワーク グループ サイズ
    • Top Memory Writes: Host and Device Global Memory
      • Buffer Address: バッファーのホスト アドレス (16 進数)
      • Context ID: ホストのコンテキスト ID
      • Command Queue ID: ホストのコマンド キュー ID
      • Start Time (ms) : 書き込み転送の開始時間 (ms)
      • Duration (ms): 書き込み転送期間 (ms)
      • Buffer Size (KB): 書き込み転送サイズ (KB)
      • Writing Rate (MB/s): 書き込みレート = (バッファー サイズ) / (期間)
    • Top Memory Reads: Host and Device Global Memory
      • Buffer Address: バッファーのホスト アドレス (16 進数)
      • Context ID: ホストのコンテキスト ID
      • Command Queue ID: ホストのコマンド キュー ID
      • Start Time (ms): 読み出し転送の開始時間 (ms)
      • Duration (ms): 読み出し転送期間 (ms)
      • Buffer Size (KB): 読み出し転送サイズ (KB)
      • Reading Rate (MB/s): 読み出しレート = (バッファー サイズ) / (期間)
  2. Kernels & Compute Units
    • Kernel Execution (includes estimated device times): スケジュールおよび実行されたすべてのカーネルのプロファイル データ サマリを表示します。
      • Kernel: カーネル名
      • Number of Enqueues: カーネルがエンキューされる回数
      • Total Time (ms): すべてのエンキューのランタイム合計 (OpenCL 実行モデルで START から END まで測定)
      • Minimum Time (ms): すべてのエンキューの最小ランタイム
      • Average Time (ms): (合計時間) / (エンキュー数)
      • Maximum Time (ms): すべてのエンキューの最大ランタイム
    • Compute Unit Utilization (includes estimated device times): FPGA のすべての計算ユニットのサマリ プロファイル データを表示します。
      • Device: デバイス名 (フォーマット: <device>-<ID>)
      • Compute Unit: 計算ユニット名
      • Kernel: 計算ユニットが関連付けられるカーネル
      • Global Work Size: カーネルの NDRange (フォーマットは x:y:z)
      • Local Work Size: ローカル ワーク グループ サイズ (フォーマットは x:y:z)
      • Number of Calls: 計算ユニットが呼び出される回数
      • Total Time (ms): すべての呼び出しのランタイム合計
      • Minimum Time (ms): すべての呼び出しの最小ランタイム
      • Average Time (ms): (合計時間) / (ワーク グループ数)
      • Maximum Time (ms): すべての呼び出しの最大ランタイム
      • Clock Frequency (MHz): アクセラレータに使用されるクロック周波数 (MHz)
  3. Data Transfers
    • Data Transfer: Host and Global Memory: ホストとデバイス メモリ間の PCI Express® リンクを介したすべての読み出しおよび書き込み転送のプロファイル データを表示します。
      • Context:Number of Devices: コンテキスト ID およびコンテキスト内のデバイス数
      • Transfer Type: [READ] または [WRITE]
      • Number of Transfers: ホスト データ転送数
        注記: printf 転送が含まれる場合があります。
      • Transfer Rate (MB/s): (送信バイト数合計) / (合計時間 (µs))

        合計時間にはソフトウェア オーバーヘッドが含まれます。

      • Average Bandwidth Utilization (%): (転送レート) / (最大転送レート)

        最大転送レート = (256/8 バイト) * (300 MHz) = 9.6 GB/s

      • Average Size (KB): (送信された合計 KB) / (転送数)
      • Total Time (ms): 転送の合計時間
      • Average Time (ms): (合計時間) / (転送数)
    • Data Transfer: Kernels and Global Memory: FPGA とデバイス メモリ間のすべての読み出しおよび書き込み転送のプロファイル データを表示します。
      • Device: デバイス名
      • Compute Unit/Port Name: <計算ユニット名>/<ポート名>
      • Kernel Arguments: このポートに接続される引数のリスト
      • DDR Bank: このポートが接続される DDR バンク数
      • Transfer Type: [READ] または [WRITE]
      • Number of Transfers: デバイスで監視される AXI トランザクションの数
        注記: printf 転送が含まれる場合があります。
      • Transfer Rate (MB/s): (送信されたバイト数合計) / (計算ユニット合計時間)
        • 計算ユニットの合計時間 = 計算ユニットの合計実行時間
        • 送信バイト数合計 = すべてのトランザクションのバイト数合計
      • Average Bandwidth Utilization (%): (転送レート) / (0.6 * 最大転送レート)

        最大転送レート = (512/8 バイト) * (300 MHz) = 19200 MB/s

      • Average Size (KB): (送信された合計 KB) / (AXI トランザクション数)
      • Average Latency (µs): (全トランザクションのレイテンシ合計) / (AXI トランザクション数)
  4. [OpenCL API Calls]: ホスト アプリケーションで実行されるすべての OpenCL ホスト API 関数呼び出しのプロファイル データを表示します。
    • API Name: API 関数の名前 (例: clCreateProgramWithBinaryclEnqueueNDRangeKernel)
    • Number of Calls: この API の呼び出し回数
    • Total Time (ms): すべての呼び出しのランタイム合計
    • Minimum Time (ms): すべての呼び出しの最小ランタイム
    • Average Time (ms): (合計時間) / (呼び出し回数)
    • Maximum Time (ms): すべての呼び出しの最大ランタイム
  5. Kernel Internals
    • Compute Units: Running Time and Stalls: 計算ユニットの実行時間 (マイクロ秒) と、その実行時間の停止時間の割合 (%) をレポートします。
      ヒント: Kernel Internals タブでは、時間をマイクロ秒 (µs) でレポートしますが、[Profile Summary] レポートのほかの部分の時間単位はミリ秒 (ms) です。
      • Compute Unit: 計算ユニットのインスタンス名を示します。
      • Running Time (µs): CU の実行時間の合計をレポートします。
      • Intra-Kernel Stream Stalls (%): カーネル間のデータ ストリーミング中に停止した実行時間の割合 (%) をレポートします。
      • External Memory Stalls (%): CU 外のメモリ転送で停止した実行時間の割合 (%) をレポートします。
      • External Stream Stalls (%): CU 内外のデータ ストリーミング中に停止した実行時間の割合 (%) をレポートします。
    • Functions: Running Time and Stalls: CU 内の関数の実行時間 (マイクロ秒) と、その実行時間の停止時間の割合 (%) をレポートします。
      • Compute Unit: 計算ユニットのインスタンス名を示します。
      • Function: CU 内の関数名を示します。
      • Running Time (µs): 関数の実行時間の合計をレポートします。
      • Intra-Kernel Stream Stalls (%): カーネル間のデータ ストリーミング中に停止した実行時間の割合 (%) をレポートします。
      • External Memory Stalls (%): CU 外のメモリ転送で停止した実行時間の割合 (%) をレポートします。
      • External Stream Stalls (%): CU 内外のデータ ストリーミング中に停止した実行時間の割合 (%) をレポートします。
    • Compute Units: Port Data Transfer: 計算ユニットの特定ポートのデータ転送をレポートします。
      • Compute Unit: 計算ユニットのインスタンス名を示します。
      • Port: 計算ユニットのポート名を示します。
      • Write Time (µs): ポートのデータ書き込み時間合計を指定します。
      • Outstanding Write (%): 書き込み処理にかかった実行時間の割合 (%) を指定します。
      • Read Time (µs): ポートのデータ読み出し時間合計を指定します。
      • Outstanding Read (%): 読み出し処理にかかった実行時間の割合 (%) を指定します。

アプリケーション タイムライン

アプリケーション タイムラインは、ホストとカーネルのイベント情報を収集し、共通のタイムラインに表示します。これは、システムの全体的な状態とパフォーマンスを視覚的に表示して理解するのに役立ちます。このグラフィカル表示により、カーネル同期および並列実行の効率に関する問題を確認できます。表示されるイベントには、次のものがあります。

  • ホスト コードからの OpenCL API 呼び出し。
  • 計算ユニット、AXI トランザクションの開始/停止を含むデバイス トレース データ。
  • ホスト イベントおよびカーネルの開始/停止。

これはアプリケーションのデバッグおよびプロファイリングには有益ですが、アプリケーション実行に余計な時間がかかるので、タイムラインとデバイス トレース データはデフォルトでは収集されません。ただし、トレース データはカーネルの専用リソースに収集されるので、カーネルの機能には影響しません。データは実行の最後にのみ解放されます (v++ --trace_memory オプション)。

次の図に、ホストおよびデバイスのイベントを共通のタイムラインに表示する [Application Timeline] ウィンドウの例を示します。ホスト アクティビティが上部に、カーネル アクティビティが下部に表示されます。ホスト アクティビティには、プログラムの作成、カーネルの実行、およびグローバル メモリとホスト間のデータ転送が含まれます。カーネル アクティビティには、読み出し/書き込みアクセス、およびグローバル メモリとカーネル間の転送が含まれます。この情報は、アプリケーション実行の詳細を理解し、パフォーマンスを向上できる部分を特定するのに有益です。

6: アプリケーション タイムライン

コマンド ライン フローでもタイムライン データの収集をイネーブルにできますが、表示には Vitis アナライザーを使用する必要があります。詳細は、Vitis アナライザーの使用 を参照してください。

アプリケーション タイムラインの生成と表示

アプリケーション タイムライン レポートを生成するには、次の手順に従って、コマンド ライン フローでタイムラインとデバイス トレース データ収集イネーブルにします。

  1. リンク時に v++ --profile_kernel オプションを使用して、カーネルにアクセラレーション モニターおよび AXI パフォーマンス モニターを追加することにより FPGA バイナリをインストルメント化します。このオプションには、Vitis コンパイラ コマンド に説明されているように、datastall、および exec の 3 つのインストルメンテーション オプションがあります。次に、v++ リンク コマンドに --profile_kernel を追加する例を示します。
    v++ -g -l --profile_kernel data:all:all:all ...
  2. ビルド プロセスでカーネルをインストルメント化したら、xrt.ini ファイルを編集してアプリケーションのランタイム実行時にデータ収集をイネーブルにする必要があります。詳細は、xrt.ini ファイル を参照してください。
    次の xrt.ini ファイルでは、アプリケーションの実行時に最大限の情報が収集されます。
    [Debug]
    profile=true
    timeline_trace=true
    data_transfer_trace=coarse
    stall_trace=all
    
    ヒント: 大量のトレース データを収集する場合、v++ コマンドで --trace_memory を指定し、xrt.initrace_buffer_size キーワードを含める必要がある場合があります。

    アプリケーションの実行後、アプリケーション タイムライン データは timeline_trace.csv という CSV ファイルに保存されます。

  3. CSV レポートは、スプレッドシート ツールまたはユーティリティ、あるいは Vitis アナライザー ツールで開くことができ、アプリケーションをビルドしたときに Vitis コンパイラから、アプリケーションを実行したときに XRT ライブラリから表示されるようになっています。Vitis アナライザーを起動してレポートを開くには、次のコマンドを使用できます。
    vitis_analyzer timeline_trace.csv

アプリケーション タイムラインの解釈

[Application Timeline] ビューは、ホストおよびデバイスのイベントを共通のタイムラインに表示します。この情報は、アプリケーション実行の詳細を理解し、パフォーマンスを向上できる部分を特定するのに有益です。アプリケーション タイムライン レポートには、[Host] と [Device] の 2 つのセクションがあります。[Host] セクションには、ホスト側から開始されるアクティビティのトレースが表示され、[Device] セクションには、FPGA の CU のアクティビティが表示されます。

レポートは、次のような構成になっています。

  • Host
    OpenCL API Calls
    すべての OpenCL API 呼び出しがトレースされます。アクティビティ時間はホストの視点から測定されます。
    General
    clCreateProgramWithBinaryclCreateContextclCreateCommandQueue などの一般的な OpenCL API 呼び出しがトレースされます。
    Queue
    特定のコマンド キューに関連する OpenCL API 呼び出しがトレースされます。これには、clEnqueueNDRangeKernelclEnqueueMigrateMemObjects などのコマンドが含まれます。ユーザー アプリケーションで複数のコマンド キューが作成された場合は、このセクションにすべてのキューとそのアクティビティが表示されます。
    Data Transfer
    ホストからデバイス メモリまでの DMA 転送がトレースされます。OpenCL ランタイムにインプリメントされる DMA スレッドは複数あり、通常は同数の DMA チャネルがあります。DMA 転送は、ユーザー アプリケーションが clEnqueueMigrateMemObjects などの OpenCL API を呼び出すことにより開始されます。これらの DMA 要求がランタイムに転送され、スレッドの 1 つに割り当てられます。ホストからデバイスまでのデータ転送は Write の下、デバイスからホストまでのデータ転送は Read の下に表示されます。
    Kernel Enqueues
    ホスト プログラムによりエンキューされたカーネルが表示されます。ここに示されるカーネルを、デバイスのカーネル/CU と混同しないようにしてください。ここではカーネルとは NDRangeKernels および OpenCL コマンド clEnqueueNDRangeKernels および clEnqueueTask で作成されるタスクを指します。これらはホストの視点から測定された時間に対してプロットされます。複数のカーネルが同時に実行されるようにスケジュールでき、実行がスケジュールされた時点からカーネル実行の終了までがトレースされます。複数のエントリがあるのは、このためです。行数は、オーバーラップするカーネル実行の数によって異なります。
    注記: 実際にはプロセスが即座に実行できない場合もあるので、カーネルのオーバーラップはデバイス上の並列実行とは異なります。
  • Device "name"
    Binary Container "name"
    バイナリ コンテナーの名前。
    Accelerator "name"
    FPGA 上の計算ユニット (アクセラレータ) の名前。
    User Functions
    Vivado HLS ツールのカーネルの場合、データフロー プロセスとしてインプリメントされる関数がトレースされます。これらの関数のトレースは、現在並列実行されているこれらの関数のアクティブなインスタンス数を示します。これらの名前は、波形がイネーブルの場合にハードウェア エミュレーションで生成されます。
    注記: 関数レベルのアクティビティは、ハードウェア エミュレーションでのみ可能です。
    • Function: "name a"
    • Function: "name b"
    Read
    CU は、AXI-MM ポートを使用して DDR から読み出しを実行します。CU で読み出されるデータのトレース データがここに表示されます。アクティビティはトランザクションとして表示され、各トランザクションのツール ヒントに詳細な AXI トランザクションが表示されます。これらの名前は --profile_kernel data が使用されると生成され、m_axi_<bundle name>(port) という形式です。
    Write
    CU は、AXI-MM ポートを使用して DDR への書き込みを実行します。CU で書き込まれるデータのトレースがここに表示されます。アクティビティはトランザクションとして表示され、各トランザクションのツール ヒントに詳細な AXI トランザクションが表示されます。これは --profile_kernel data が使用されると生成され、m_axi_<bundle name>(port) という形式です。

[Waveform] ビューおよびライブ波形ビューアー

Vitis コア開発キットでは、ハードウェア エミュレーションを実行したときに [Waveform] ビューが表示されます。[Waveform] ビューには、システム レベル、CU レベル、および関数レベルの詳細が表示されます。表示される詳細には、カーネルとグローバル メモリ間のデータ転送、カーネル パイプ間のデータフローが含まれます。これらの詳細を利用すると、システム レベルから個々の関数呼び出しまでのパフォーマンスのボトルネックを理解して、アプリケーションを最適化できるようになります。

ライブ波形ビューアーは、[Waveform] ビューに似ていますが、さらに下位レベルの詳細が表示され、ある程度のインタラクティブ機能も提供されています。ライブ波形ビューアーは、Vivado ロジック シミュレータ xsim を使用して開くこともできます。

注記: [Waveform] ビューを使用すると、Vitis アナライザー (Vitis アナライザーの使用 を参照) からデバイス トランザクションを直接確認できます。ライブ波形ビューアーは、Vivado シミュレーション波形ビューアーに開き、選択した信号に加え、ハードウェア トランザクションを確認できます。

[Waveform] ビューおよびライブ波形ビューアーのデータを収集するには、ランタイムでハードウェア エミュレーション中にシミュレーション波形を生成する必要があり、時間もディスク容量も消費するので、デフォルトでは実行されません。これらの機能をイネーブルにする方法は、波形レポートの生成と表示 を参照してください。

7: [Waveform] ビュー

または、Linux コマンド ラインから Vivado ロジック シミュレータで波形データベース (.wdb) ファイルを開きます。

xsim -gui <filename.wdb> &
ヒント: .wdb ファイルは、コンパイルされたホスト コードが実行されるディレクトリに含まれます。

波形レポートの生成と表示

ハードウェア エミュレーション中にコマンド ラインから波形データの収集をイネーブルにしてビューアーで開くには、次の手順に従います。

  1. コンパイルおよびリンクを実行するときに、-g オプションを指定してデバッグ コードの生成をイネーブルにします。
    v++ -c -g -t hw_emu ...
  2. ホスト実行ファイルと同じディレクトリに、次の次内容の xrt.ini ファイルを作成します (詳細は xrt.ini ファイル を参照)。
    [Debug]
    profile=true
    timeline_trace=true
    
    [Emulation]
    launch_waveform=batch
    ライブ波形ビューアーの場合は、launch_waveform は次のようになります。
    [Emulation]
    launch_waveform=gui
    ヒント: ライブ波形ビューアーをイネーブルにしている場合は、ハードウェア エミュレーションの実行時にシミュレーション波形が開きます。
  3. アプリケーションの実行の手順に従って、アプリケーションのハードウェア エミュレーション ビルドを実行します。ハードウェア トランザクション データは、<hardware_platform>-<device_id>-<xclbin_name>.wdb 波形データベース ファイルに収集されます。このファイルの場所は、ディレクトリ構造 を参照してください。
  4. [Waveform] ビューおよびライブ波形ビューアー の説明に従って、Vitis で [Waveform] ビューを開きます。

[Waveform] ビューのデータの解釈

次の図に、[Waveform] ビューを示します。

8: [Waveform] ビュー

[Waveform] ビューおよびライブ波形ビューアーは、ナビゲートしやすいように階層形式で表示されます。

  • [Waveform] ビューはハードウェア エミュレーション (カーネル トレース) 中に実際に生成された波形に基づいているので、抽象化されたデータの基になっている個別信号まで詳細に表示できます。ただし、[Waveform] ビューは後処理されたデータから生成されるので、信号を追加することはできず、DATAFLOW トランザクションなどのランタイム解析の一部は表示されません。
  • ライブ波形ビューアーは Vivado ロジック シミュレータ (xsim) の実行の一部として表示されるので、RTL デザインの信号および内部信号をライブ波形に追加できます。波形ビューアーの詳細は、『Vivado Design Suite ユーザー ガイド: ロジック シミュレーション』 (UG900) を参照してください。

[Waveform] ビューおよびライブ波形ビューアーには、次のような階層があります。

Device "name"
ターゲット デバイスの名前。
Binary Container "name"
バイナリ コンテナーの名前。
Memory Data Transfers
ホストからバンクに到着する読み出しおよび書き込みトランザクションすべてのトレースを DDR バンクごとに示します。
Kernel "name" 1:1:1
各カーネルおよびそのカーネルの計算ユニットごとに、計算ユニットからのアクティビティを示します。
Compute Unit: "name"
計算ユニットの名前。
CU Stalls (%)
外部メモリ アクセス、内部ストリーム (データフロー)、または外部ストリーム (OpenCL パイプ) が原因で回路の一部がストールした場合に通知するため、Vivado HLS ツールからストール信号が提供されています。カーネル トレースに詳細に示されるストール バスは、最下位のストール信号をすべてコンパイルし、すべての時点でのストールしている割合 (%) をレポートします。これにより、カーネルのどの程度がシミュレーションでストールしているかがわかります。

たとえば、100 個の最下位ストール信号があり、あるクロック サイクルでそのうち 10 個がアクティブであれば、[CU Stall (%)] は 10% です。そのうちの 1 つがアクティブでなくなれば、9% になります。

Data Transfers
計算ユニットの各マスター AXI ポートから DDR への読み出し/書き込みデータ転送アクセスを示します。
User Functions
HLS カーネル用に表示される情報で、ユーザー関数を示します。
Function: "name"
関数名。
Dataflow/Pipeline Activity
関数がデータフロー プロセスとしてインプリメントされる場合、並列実行される関数の数を示します。
Active Iterations
現在のアクティブなデータフローのイテレーションを示します。すべての同時実行を表示するため、行数は動的に増加します。
StallNoContinue
データフロー プロセスにより出力でストールが発生した (関数は完了したが隣接するデータフロー プロセスからの続行信号は受信していない) ことを示すストール信号です。
RTL Signals
データフロー プロセスの上記のトランザクション ビューを解釈するために使用された RTL 制御信号です。
Function Stalls
プロセスで発生したさまざまなタイプのストールを示します。
External Memory
DDR メモリにアクセス中に発生したストール。
Internal-Kernel Pipe
計算ユニットがパイプを使用して通信する場合に、関連するストールを示します。
Intra-Kernel Dataflow
内部からカーネルへの FIFO アクティビティ。
Function I/O
実際のインターフェイス信号。
Function: "name"
関数名。
Function: "name"
関数名。