Profiling the Application

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. This feature provides two types of custom profiling:

User range

Profiles the specified start/end times across a range of code. This captures the span of time within which an action occurs in the host application.

User events

Marks an event in the timeline. The user event is added to the timeline waveform at whatever point in time it occurs.

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 xrt/core/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 summary, 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. 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]
   profile = 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 profile_summary.csv report file when running the application, and also creates the profile_kernels.csv and timeline_kernels.csv files 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 Application Timeline 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 Application Timeline 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 will show the related stalls.

Intra-Kernel Dataflow

FIFO activity internal to the kernel.

Function I/O

Actual interface signals.

Function: "name"

Function name.

Function: "name"

Function name.