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

ホスト アプリケーションからのすべての XRT 関連動作は、OpenCL API 呼び出しまたは XRT API 呼び出しのいずれかを使用して、プロファイリングのために自動的にトラックされます。ただし、XRT 関連イベント以外のホスト アプリケーションをプロファイルして、ユーザー指定の動作またはイベントに基づいてイベント データを取得することもできます。この機能には、次の 2 種類のカスタム プロファイリングがあります。

[User range]

指定した開始時刻と終了時刻をコード範囲全体でプロファイルします。これにより、ホスト アプリケーションで動作が発生する期間が取得されます。

[User events]

タイムライン内のイベントをマークします。ユーザー イベントは、発生した時点でタイムライン波形に追加されます。

カスタム プロファイリングを使用するには、ホスト アプリケーションのソース コードとビルド プロセスをいくつか変更する必要があります。次に説明するように、コード内で C または C++ API を使用する必要があります。また、ホスト アプリケーションをリンクするときに、xrt_coreutil ライブラリを含める必要もあります。

• C/C++ API については次に説明しますが、https://github.com/Xilinx/XRT/blob/master/src/runtime_src/core/include/experimental/xrt_profile.hからも参照できます。
• C と C++ の両方に次を追加する必要があります。

  #include xrt/core/include/experimental/xrt_profile.h

• ホスト コードをリンクする場合は、コンパイラのコマンドラインに -lxrt_coreutil を追加します。

Vitis ソフトウェア プラットフォームでは、実行時間にほとんど影響のない、最小限の情報を提供するオーバーヘッドの少ないプロファイリングがサポートされています。ランタイム中にこのオプションを使用した場合、タイムライン トレースは表示されますが、情報量は限られたものになります。オーバーヘッドの少ないプロファイリングは、OpenCL イベントで最小限の情報を取り込み、実行の終了時に lop_trace.csv という CSV ファイルをダンプします。オーバーヘッドが少ないプロファイリングは、3 つのフロー (ハードウェア、ハードウェア エミュレーション、およびソフトウェア エミュレーション) すべてで実行できます。

オーバーヘッドの少ないプロファイリングをイネーブルにするには、xrt.ini ファイル の [Debug] セクションに示すように lop_trace という新しいオプションを使用します。デフォルトでは、lop_trace は FALSE なので、イネーブルにする場合は ini パラメーターを TRUE に設定する必要があります。

xrt.ini file
[Debug]
lop_trace=true

lop_trace=true をイネーブルにすると、ランタイムが Vitis アナライザーの Run Summary で表示可能な lop_trace.csv ファイルが生成されます。

vitis_analyzer <project>.run_summary

最小限のオーバーヘッドを取得するため、標準の OpenCL プロファイリングで収集される情報は省かれます。特に、オーバーヘッドの少ないプロファイリング トレースでは次の情報含まれません。

• 計算ユニット実行またはカーネル メモリ転送などのデバイス イベント
• デスティネーション アドレスまたはサイズなどのメモリ読み出しまたは書き込みに関する情報
• カーネル名や NDRange サイズなどのカーネル エンキューに関する情報
• バッファー転送とカーネル エンキュー間の依存

カーネルがコンパイルされたとき、および 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

ルールを満たすために変更可能なモデル ソース コードまたはツール変更によく使用される方法へのリンクが含まれます。リンクをクリックすると、特定の問題に適用可能なヒントやコードを含むポップアップ ウィンドウまたは資料が開きます。

システム見積もりレポートは、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
----------------------------------------------------------------------------

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

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

vitis_analyzer <output_filename>.compile_summary

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

HLS 合成レポートは、左側の列にモジュール階層を示すスプレッドシートです。HLS の実行により生成されたモジュールおよびループが、この階層に表示されます。HLS 合成レポートには、次の情報が含まれます。

• 違反のタイプ
• クロック サイクル数で示されたレイテンシ
• 絶対時間 (µs) で示されたレイテンシ
• 反復レイテンシ
• 反復間隔
• ループのトリップカウント
• パイプライン処理の有無
• BRAM、DSP、FF、および LUT の使用率見積もり
• 負のスラック

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

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

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

   [Debug]
   profile = true

3. [Kernel Internals] データのプロファイリングをイネーブルにするには、xrt.ini の [Emulation] セクションに debug_mode タグも追加する必要があります。

   [Emulation]
   debug_mode = batch

デバイス バイナリおよび xrt.ini ファイルでプロファイリングをイネーブルにすると、アプリケーションを実行したときに、ランタイムでprofile_summary.csv レポート ファイルが生成されるほか、Kernel Internals をオンにしたときに profile_kernels.csv および timeline_kernels.csv ファイルも作成されます。これらのファイルは Profile Summary レポートにリンクされ、Run Summary を使用して Vitis アナライザー ツールで表示できます。次のコマンドで Run Summary を開きます。

vitis_analyzer <project>.run_summary

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

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

• Host

  OpenCL API Calls

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

  General

  clCreateProgramWithBinary、clCreateContext、clCreateCommandQueue などの一般的な OpenCL API 呼び出しがトレースされます。

  Queue

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

  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

  Vitis HLS ツールのカーネルの場合、データフロー プロセスとしてインプリメントされる関数がトレースされます。これらの関数のトレースは、現在並列実行されているこれらの関数のアクティブなインスタンス数を示します。これらの名前は、波形がイネーブルの場合にハードウェア エミュレーションで生成されます。

  注記: 関数レベルのアクティビティは、ハードウェア エミュレーションでのみ可能です。

  • Function: "name a"
  • Function: "name b"

  Read

  CU は、AXI-MM ポートを使用して DDR から読み出しを実行します。CU で読み出されるデータのトレース データがここに表示されます。アクティビティはトランザクションとして表示され、各トランザクションのツール ヒントに詳細な AXI トランザクションが表示されます。これらの名前は、--profile.data がを CU のために使用されると生成されます。

  Write

  CU は、AXI-MM ポートを使用して DDR への書き込みを実行します。CU で書き込まれるデータのトレースがここに表示されます。アクティビティはトランザクションとして表示され、各トランザクションのツール ヒントに詳細な AXI トランザクションが表示されます。これは、--profile.data が CU 用に指定されたときに生成されます。

次の図に、[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 パイプ) が原因で回路の一部がストールした場合に通知するため、Vitis 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"

関数名。