Profiling the Application

The Vitis tool supports recording continuous trace data while the application is running. The application can run for a very long time thus leading to the capture of significant trace data, which can result in issues like incomplete trace data especially when the memory resource used for trace data is not large enough. Using continuous trace, analysis of the trace can be carried out while the application is still running or if the application has crashed before completion.

With the ability to continuously capture trace data, the Timeline Trace and Application Timeline reports can be dynamically updated in the Vitis analyzer tool while your application is running. Once these reports are loaded in Vitis Analyzer, there is a hyperlink available indicating that the current report is being modified on the disk. If new data needs to be loaded, Reload or Auto-Reload options are available on the banner to let you view the updated report as your application runs and trace data is generated.

Continuous trace is not enabled by default. Additionally, the memory resources of an FPGA are not unlimited. So if the application generates large trace data, a circular buffer for storing the data can be used. The circular buffer can be written, offloaded to the host, and reused again. By enabling a circular buffer with continuous trace, the memory resources needed are even smaller thus saving available resources on the device. However, an application run with continuous trace/circular buffer may result in multiple device trace files.

Here are some scenarios where it is recommended to use the memory resource as a circular buffer.

The circular buffer implementation is automatically turned on when continuous trace is enabled in the xrt.ini. The flow requires the following settings for enabling continuous trace.

• In the xrt.ini file, continuous_trace is set to TRUE
• v++ linking option --trace_memory is set to DDR or HBM

You can optionally set:

• The size of the trace buffer using trace_buffer_size in the xrt.ini file. This defaults to 1 MB.
• The interval at which the trace buffer is offloaded from the device using trace_buffer_offload_interval_ms in the xrt.ini file. The default is 10 ms.
• The interval at which files are dumped by setting trace_file_dump_interval_s. The default is 3 seconds.

[XRT] WARNING: Unable to use circular buffer for continuous trace offload. Please increase trace buffer size and/or reduce continuous
trace interval. Minimum required offload rate (bytes per second) : 104857600 Requested offload rate : 819200

[Debug]
opencl_summary=true
opencl_trace=true
data_transfer_trace=coarse
stall_trace=all
continuous_trace=true
// The following are optional and needed only in rare circumstances

trace_buffer_size=20M
trace_buffer_offload_interval_ms=10
trace_file_dump_interval_s=2

The following are the results of these settings:

• opencl_summary: Enables the generation of host-related OpenCL API profile summary report, opencl_summary.csv file is created.
• opencl_trace: Enables the generation of host-related OpenCL API trace, opencl_trace.csv files is created.
• data_transfer_trace: Enables the collection of kernel activity to be added to profile summary and trace, device_trace_0.csv files are created with 0 being the device number.
• stall_trace: Enables the hardware generation of stalls into compute units.
• continuous_trace: Enables the continuous dumping of files for trace and the continuous reading of device data into the host.
• trace_buffer_offload_interval_ms: Controls the reading of device data from the device to the host in milliseconds.
• trace_file_dump_interval_s: Controls the time between dumping of trace files in seconds.

As a result, there are several CSV files generated in addition to the xclbin.run_summary as part of the application run using the above xrt.ini file. Vitis Analyzer only needs the generated run_summary file and will use the relevant CSV files to display the profile summary and timeline trace.

Here are the recommendations on setting up an application for trace data dumping:

1. By default, an 8k FIFO is used for saving trace data. The FIFO size can be increased but not preferred above 64k and needs to be preallocated as part of the v++ linking step. It is also preferred to use device memory for saving trace data. If you specify a memory bank for trace, you can use trace_buffer_size option in xrt.ini to control the amount of trace generated at runtime. With device memory, the default size is 1M and the maximum size is 4095M.
2. If still unable to dump maximum trace, disable stall trace by setting stall_trace=off or stall_trace=on with data_transfer_trace=coarse.
3. If the application requires larger size of trace buffer, enable circular buffer by setting continuous_trace=true with default settings of trace_buffer_offload_interval_ms=10 and trace_file_dump_interval_s=5. Ideally, a continuous trace feature should be used for the following cases:
   • Long-running design with minimal trace generated
   • Debugging application crashes where some .csv files might still be available for debugging
4. If the application run is still unable to dump the maximum trace, the trace_buffer_size can further be increased.
5. If the application still creates huge trace data that the host cannot keep up, use the smaller size of trace_file_dump_interval, which creates multiple files equivalent to the interval provided.
6. Lastly, continuous trace can generate several trace files as part of the application run in addition to xclbin.run_summary file. The Vitis Analyzer only needs the generated run_summary file and can pick the relevant CSV files generated to display profile summary and timeline trace to provide a better experience.

All XRT related actions from the host application are automatically tracked for profiling, through either theOpenCL API calls, or the XRT API calls. However, you can also profile the host application beyond the XRT related events, capturing event data based on user-specified actions or events.

The user_range and user_event data can be captured to the Profile Summary and Timeline Trace reports for display in Vitis analyzer. As seen in the figure below, the Profile Summary shows the number of occurrences of a given event and the range. The User Ranges table also reports the Min/Max/Avg/Total duration of the user-defined ranges in the host code. In the Timeline Trace report user_range elements in the host code are displayed in a separate row, and user_event markers are added at specific points on the timeline.

Using custom profiling requires a few changes in your host application source code and build process. You must make use of C or C++ API in your code, as described below, and you must include the xrt_coreutil library when linking your host application.

• The C/C++ API are described below, but can also be found at the following URL: https://github.com/Xilinx/XRT/blob/master/src/runtime_src/core/include/experimental/xrt_profile.h.
• For both C and C++ you must add the following:

  #include experimental/xrt_profile.h

• When linking host code, add -lxrt_coreutil to the compiler command line.

The Vitis software platform supports low overhead profiling that provides minimal information with little effect on execution time. Using this option during runtime, the timeline trace is still available but with a reduced amount of information. Low overhead profiling captures minimal information on OpenCL events and dumps a CSV file called lop_trace.csv at the end of execution. Low overhead profiling can be run in all three flows (hardware, hardware emulation, and software emulation).

To enable low overhead profiling, there is a new flag in the "Debug" section of the xrt.ini File called lop_trace. By default, lop_trace is FALSE and must be enabled by setting the ini parameter to TRUE.

xrt.ini file
[Debug]
lop_trace=true

When lop_trace=true is enabled, the runtime will generate lop_trace.csv which can be viewed in the Run Summary within Vitis analyzer.

vitis_analyzer <project>.run_summary

To obtain the lowest possible overhead, information collected in normal OpenCL profiling is omitted. Specifically, the following information is expected to not be available in the low overhead profiling trace:

• Device events, such as compute unit executions or kernel memory transfers
• Information about memory reads or writes, such as destination address or size
• Information about kernel enqueues, such as kernel name or NDRange sizes
• Dependencies between buffer transfers and kernel enqueue

When kernels are compiled and when the FPGA binary is linked, guidance reports are generated automatically by the v++ command. You can view these reports in the Vitis analyzer by opening the <output_filename>.compile_summary or the <output_filename>.link_summary for the application project. The <output_filename> is the output of the v++ command.

As an example, launch the Vitis analyzer and open the report using this command:

vitis_analyzer <output_filename>.link_summary

When the Vitis analyzer opens, it displays the link summary report, as well as the compile summaries, and a collection of reports generated during the compile and link processes. Both the compile and link steps generate Guidance reports to view by clicking the Build heading on the left-hand side. Refer to Using the Vitis Analyzer for more information.

The Guidance view places each entry in a separate row. Each row might contain the name of the guidance rule, threshold value, actual value, and a brief but specific description of the rule. The last field provides a link to reference material intended to assist in understanding and resolving any of the rule violations.

In the GUI Guidance view, guidance rules are grouped by categories and unique IDs in the Name column and annotated with symbols representing the severity. These are listed individually in the HTML report. In addition, as the HTML report does not show tooltips, a full Name column is included in the HTML report as well.

The following list describes all fields and their purpose as included in the HTML guidance reports.

Id

Each guidance rule is assigned a unique ID. Use this id to uniquely identify a specific message from the guidance report.

Name

The Name column displays a mnemonic name uniquely identifying the guidance rule. These names are designed to assist in memorizing specific guidance rules in the view.

Severity

The Severity column allows the easy identification of the importance of a guidance rule.

Full Name

The Full Name provides a less cryptic name compared to the mnemonic name in the Name column.

Categories

Most messages are grouped within different categories. This allows the GUI to display groups of messages within logical categories under common tree nodes in the Guidance view.

Threshold

The Threshold column displays an expected threshold value, which determines whether or not a rule is met. The threshold values are determined from many applications that follow good design and coding practices.

Actual

The Actual column displays the values actually encountered on the specific design. This value is compared against the expected value to see if the rule is met.

Details

The Details column provides a brief message describing the specifics of the current rule.

Resolution

The Resolution column provides a pointer to common ways the model source code or tool transformations can be modified to meet the current rule. Clicking the link brings up a popup window or the documentation with tips and code snippets that you can apply to the specific issue.

The System Estimate report can be opened in the Vitis analyzer tool, intended for viewing reports from the Vitis compiler when the application is built, and the XRT library when the application is run. You can launch the Vitis analyzer and open the report using the following command:

vitis_analyzer <output_filename>.link_summary

The <output_filename> is the output of the v++ command. This opens the Link Summary for the application project in the Vitis analyzer tool. Then, select the System Estimate report. Refer to Using the Vitis Analyzer for more information.

The System Estimate report generated by the v++ command provides information on every binary container in the application, as well as every compute unit in the design. The report is structured as follows:

• Target device information
• Summary of every kernel in the application
• Detailed information on every binary container in the solution

The following example report file represents the information generated for the estimate report:

----------------------------------------------------------------------------
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
----------------------------------------------------------------------------

The HLS report can be viewed through the Vitis analyzer by opening the <output_filename>.compile_summary or the <output_filename>.link_summary for the application project. The <output_filename> is the output of the v++ command.

You can launch the Vitis analyzer and open the report using the following command:

vitis_analyzer <output_filename>.compile_summary

When the Vitis analyzer opens, it displays the Compile Summary and a collection of reports generated during the compile process. Refer to Using the Vitis Analyzer for more information.

The HLS Synthesis report is a spreadsheet listing the module hierarchy in the left column. This section is describing one section of the HLS report: Performance and Resource Estimates. Each module and loop generated by the HLS run is represented in this hierarchy. The HLS Synthesis report contains the following columns:

• Violation Type
• Latency in clock cycles
• Latency in absolute time (µs)
• Iteration latency
• Iteration Interval
• Loop Tripcount
• Pipelined
• Utilization Estimates of BRAM, DSP, FF, and LUT
• Negative Slack

If this information is part of a hierarchical block, it will sum up the information of the blocks contained in the hierarchy. Therefore, the hierarchy can also be navigated from within the report when it is clear which instance contributes to the overall design.

Capturing the data required for the Profile Summary requires a few steps prior to actually running the application.

1. The FPGA binary (xclbin) file is configured for capturing profiling data by default. However, using the v++ --profile option during the linking process enables a greater level of detail in the profiling data captured. For more information, see the --profile Options.
2. The runtime requires the presence of an xrt.ini file, as described in xrt.ini File, that includes the keyword for capturing profiling data:

   [Debug]
   opencl_summary = true
   opencl_device_counter=true 

3. To enable the profiling of Kernel Internals data, you must also add the debug_mode tag in the [Emulation] section of the xrt.ini:

   [Emulation]
   debug_mode = batch

With profiling enabled in the device binary and in the xrt.ini file, the runtime creates the opencl_summary.csv report file when running the application, and also creates the profile_kernels.csv file when Kernel Internals is enabled. These files are linked to the Profile Summary report which can be viewed in the Vitis analyzer tool through the Run Summary. Open the Run Summary using the following command:

vitis_analyzer <project>.run_summary

The Timeline Trace window displays host and device events on a common timeline. This information helps you understand details of application execution and identify potential areas for improvements. The Timeline Trace report has two main sections: Host and Device. The Host section shows the trace of all the activity originating from the host side. The Device section shows the activity of the CUs on the FPGA.

The report has the following structure:

• Host

  OpenCL API Calls

  All OpenCL API calls are traced here. The activity time is measured from the host perspective.

  General

  All general OpenCL API calls such as clCreateProgramWithBinary, clCreateContext, and clCreateCommandQueue, are traced here.

  Queue

  OpenCL API calls that are associated with a specific command queue are traced here. This includes commands such as clEnqueueMigrateMemObjects, and clEnqueueNDRangeKernel. If the user application creates multiple command queues, then this section shows all the queues and activities.

  Data Transfer

  In this section the DMA transfers from the host to the device memory are traced. There are multiple DMA threads implemented in the OpenCL runtime and there is typically an equal number of DMA channels. The DMA transfer is initiated by the user application by calling OpenCL APIs such as clEnqueueMigrateMemObjects. These DMA requests are forwarded to the runtime which delegates to one of the threads. The data transfer from the host to the device appear under Write as they are written by the host, and the transfers from device to host appear under Read.

  Kernel Enqueues

  The kernels enqueued by the host program are shown here. The kernels here should not be confused with the kernels/CUs on the device. Here kernel refers to the NDRangeKernels and tasks created by the OpenCL commands clEnqueueNDRangeKernels and clEnqueueTask. These are plotted against the time measured from the host’s perspective. Multiple kernels can be scheduled to be executed at the same time, and they are traced from the point they are scheduled to run until the end of the kernel execution. This is the reason for multiple entries. The number of rows depend on the number of overlapping kernel executions.

  Note: Overlapping of the kernels should not be mistaken for actual parallel execution on the device as the process might not be ready to execute right away.

• Device "name"

  Binary Container "name"

  Binary container name.

  Accelerator "name"

  Name of the compute unit (a.k.a., Accelerator) on the FPGA.

  User Functions

  In the case of the Vitis HLS tool kernels, functions that are implemented as data flow processes are traced here. The trace for these functions show the number of active instances of these functions that are currently executing in parallel. These names are generated in hardware emulation when waveform is enabled.

  Note: Function level activity is only possible in hardware emulation.

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

  Read

  A CU reads from the DDR over AXI-MM ports. The trace of a data read by a CU is shown here. The activity is shown as transaction and the tool-tip for each transaction shows more details of the AXI transaction. These names are generated when --profile.data is for the CU.

  Write

  A CU writes to the DDR over AXI-MM ports. The trace of data written by a CU is shown here. The activity is shown as transactions and the tool-tip for each transaction shows more details of the AXI transaction. This is generated when --profile.data is specified for the CU.

The following image shows the Waveform view:

The Waveform and Live Waveform views are organized hierarchically for easy navigation.

• The Waveform view is based on the actual waveforms generated during hardware emulation (Kernel Trace). This allows the viewer to descend all the way down to the individual signals responsible for the abstracted data. However, because the Waveform view is generated from the post-processed data, no additional signals can be added to the report, and some of the runtime analysis cannot be visualized, such as DATAFLOW transactions.
• The Live Waveform viewer is displaying the Vivado logic simulator (xsim) run, so you can add extra signals and internals of the register transfer (RTL) design to the live view. Refer to the Vivado Design Suite User Guide: Logic Simulation (UG900) for information on working with the Waveform viewer.

The hierarchy of the Waveform and Live Waveform views include the following:

Device "name"

Target device name.

Binary Container "name"

Binary container name.

Memory Data Transfers

For each DDR Bank, this shows the trace of all the read and write request transactions arriving at the bank from the host.

Kernel "name" 1:1:1

For each kernel and for each compute unit of that kernel, this section breaks down the activities originating from the compute unit.

Compute Unit: "name"

Compute unit name.

CU Stalls (%)

Stall signals are provided by the Vitis HLS tool to inform you when a portion of the circuit is stalling because of external memory accesses, internal streams (that is, dataflow), or external streams (that is, OpenCL pipes). The stall bus shown in detailed kernel trace compiles all of the lowest level stall signals and reports the percentage that are stalling at any point in time. This provides a factor of how much of the kernel is stalling at any point in the simulation.

For example, if there are 100 lowest level stall signals and 10 are active on a given clock cycle, then the CU Stall percentage is 10%. If one goes inactive, then it is 9%.

Data Transfers

This shows the read/write data transfer accesses originating from each Master AXI port of the compute unit to the DDR.

User Functions

This information is available for the HLS kernels and shows the user functions.

Function: "name"

Function name.

Dataflow/Pipeline Activity

This shows the number of parallel executions of the function if the function is implemented as a dataflow process.

Active Iterations

This shows the currently active iterations of the dataflow. The number of rows is dynamically incremented to accommodate the visualization of any concurrent execution.

StallNoContinue

This is a stall signal that tells if there were any output stalls experienced by the dataflow processes (function is done, but it has not received a continue from the adjacent dataflow process).

RTL Signals

These are the underlying RTL control signals that were used to interpret the above transaction view of the dataflow process.

Function Stalls

Shows the different types of stalls experienced by the process.

External Memory

Stalls experienced while accessing the DDR memory.

Internal-Kernel Pipe

If the compute units communicated between each other through pipes, then this shows the related stalls.

Intra-Kernel Dataflow

FIFO activity internal to the kernel.

Function I/O

Actual interface signals.