Debugging Applications and Kernels

There are two kinds of commands which can be called from the gdb command line:

1. xprint commands that give visibility into XRT library data structures (cl_command_queue, cl_event, and cl_mem). These commands are explained below.
2. xstatus commands that give visibility into IP running on the Vitis target platform when debugging during hardware execution.

You can get more information about the xprint and xstatus commands by using the help <command> from the gdb command prompt.

A typical application for these commands is when you see the host application hang. In this case, the host application could be waiting for the command queue to finish, or waiting on an event list. Printing the command queue using the xprint queue command can tell you what events are unfinished, allowing you to analyze dependencies between events.

The output of both of these commands is automatically tracked when debugging with the Vitis IDE. In this case, three tabs are provided next to the common tabs for Variables, Breakpoints, and Registers in the upper left corner of the debug perspective. These are labeled Command Queue, Memory Buffers, and Platform Debug, showing the output of xprint queue, xprint mem, and xstatus, respectively.

GDB kernel debugging is supported for the software emulation and hardware emulation flows. When the GDB executable is connected to the kernel in the IDE or command line flows, you can set breakpoints and query the content of variables in the kernel, similar to normal host code debugging. This is fully supported in the software emulation flow because the kernel GDB processes attach to the spawned software processes.

However, during hardware emulation, the kernel source code is transformed into RTL, created by Vitis HLS, and executed. As the RTL model is simulated, all transformations for performance optimization and concurrent hardware execution are applied. For that reason, not all C/C++/OpenCL lines can be uniquely mapped to the RTL code, and only limited breakpoints are supported and at only specific variables can be queried. Today, the GDB tool therefore breaks on the next possible line based on requested breakpoint statements and clearly states if variables can not be queried based on the RTL transformations.

For OpenCL kernels, additional runtime checks can be performed during software emulation. These additional checks include:

• Checking whether an OpenCL kernel makes out-of-bounds accesses to the interface buffers (fsanitize=address).
• Checking whether the kernel makes accesses to uninitialized local memory (fsanitize=memory).

These are Vitis compiler options that are enabled through the --advanced compiler option as described in --advanced Options, using the following command syntax:

--advanced.param compiler.fsanitize=address,memory

When applied, the emulation run produces a debug log with emulation diagnostic messages that are written to <project_dir>/Emulation-SW/<proj_name>-Default>/emulation_debug.log.

The fsanitize directive can also be specified in a config file, as follows:

[advanced]
#param=<param_type>:<param_name>.<value>
param=compiler.fsanitize=address,memory

Then the config file is specified on the v++ command line:

v++ -l –t sw_emu --config ./advanced.cfg -o bin_kernel.xclbin

Refer to the Vitis Compiler Configuration File for more information on the --config option.

For C/C++ kernel models, printf() is only supported during software emulation and should be excluded from the Vitis HLS synthesis step. In this case, any printf() statement should be surrounded by the following compiler macros:

#ifndef __SYNTHESIS__
    printf("Checkpoint 1 reached");
#endif

For C++ kernels, you can also use cout in your code to add checkpoints or messages used for debugging the code. For example, you might add the following:

std::cout << "TEST " << (match ? "PASSED" : "FAILED") << std::endl;

The Xilinx Runtime (XRT) library supports the OpenCL™ printf() built-in function within kernels in all build configurations: software emulation, hardware emulation, and during hardware execution.

The following is an example of using printf() in the kernel, and the output when the kernel is executed with global size of 8:

__kernel __attribute__ ((reqd_work_group_size(1, 1, 1)))
void hello_world(__global int *a)
{
    int idx = get_global_id(0);

     printf("Hello world from work item %d\n", idx);
     a[idx] = idx;
}

The output is as follows:

Hello world from work item 0
Hello world from work item 1
Hello world from work item 2
Hello world from work item 3
Hello world from work item 4
Hello world from work item 5
Hello world from work item 6
Hello world from work item 7

You can also enable debugging in RTL kernels by manually adding ChipScope debug cores like the ILA and VIO in your RTL kernel code before packaging it for use in the Vitis development flow. From within the Vivado Design Suite, edit the RTL kernel code to manually instantiate an ILA debug core, or VIO IP from the Xilinx IP catalog, similar to using any other IP in Vivado IDE. Refer to the HDL Instantiation flow in the Vivado Design Suite User Guide: Programming and Debugging (UG908) to learn more about adding debug cores to your design.

The best time to add debug cores to your RTL kernel is when you create it. However, debug cores consume device resources and can affect performance, so it is good practice to make one kernel for debug and a second kernel for production use. The rtl_vadd_hw_debug of the RTL Kernels examples on GitHub shows an ILA debug core instantiated into the RTL kernel source file. The ILA monitors the output of the combinatorial adder as specified in the src/hdl/krnl_vadd_rtl_int.sv file.

	// ILA monitoring combinatorial adder
	ila_0 i_ila_0 (
		.clk(ap_clk),              // input wire        clk
		.probe0(areset),           // input wire [0:0]  probe0  
		.probe1(rd_fifo_tvalid_n), // input wire [0:0]  probe1 
		.probe2(rd_fifo_tready),   // input wire [0:0]  probe2 
		.probe3(rd_fifo_tdata),    // input wire [63:0] probe3 
		.probe4(adder_tvalid),     // input wire [0:0]  probe4 
		.probe5(adder_tready_n),   // input wire [0:0]  probe5 
		.probe6(adder_tdata)       // input wire [31:0] probe6
	);

You can also add the ILA debug core using a Tcl script from within an open Vivado project, using the Netlist Insertion flow described in Vivado Design Suite User Guide: Programming and Debugging (UG908), as shown in the following Tcl script example:

create_ip -name ila -vendor xilinx.com -library ip -version 6.2 -module_name ila_0
set_property -dict [list CONFIG.C_PROBE6_WIDTH {32} CONFIG.C_PROBE3_WIDTH {64} \
CONFIG.C_NUM_OF_PROBES {7} CONFIG.C_EN_STRG_QUAL {1} CONFIG.C_INPUT_PIPE_STAGES {2} \
CONFIG.C_ADV_TRIGGER {true} CONFIG.ALL_PROBE_SAME_MU_CNT {4} CONFIG.C_PROBE6_MU_CNT {4} \
CONFIG.C_PROBE5_MU_CNT {4} CONFIG.C_PROBE4_MU_CNT {4} CONFIG.C_PROBE3_MU_CNT {4} \
CONFIG.C_PROBE2_MU_CNT {4} CONFIG.C_PROBE1_MU_CNT {4} CONFIG.C_PROBE0_MU_CNT {4}] [get_ips ila_0]

After the RTL kernel has been instrumented for debug with the appropriate debug cores, you can analyze the hardware in the Vivado hardware manager as described in Debugging with ChipScope.

Supporting hardware debugging requires the platform to support several IP components, most notably the Debug Bridge. Talk to your platform designer to determine if these components are included in the target platform. If a Xilinx platform is used, debug availability can be verified using the platforminfo utility to query the platform. Debug capabilities are listed under the chipscope_debug objects.

For example, to query the a platform for hardware debug support, the following platforminfo command can be used:

$ platforminfo --json="hardwarePlatform.extensions.chipscope_debug" xilinx_u200_xdma_201830_2
{
    "debug_networks": {
        "user": {
            "name": "User Debug Network",
            "pcie_pf": "1",
            "bar_number": "0",
            "axi_baseaddr": "0x000C0000",
            "supports_jtag_fallback": "false",
            "supports_microblaze_debug": "true",
            "is_user_visible": "true"
        },
        "mgmt": {
            "name": "Management Debug Network",
            "pcie_pf": "0",
            "bar_number": "0",
            "axi_baseaddr": "0x001C0000",
            "supports_jtag_fallback": "true",
            "supports_microblaze_debug": "true",
            "is_user_visible": "false"
        }
    }
}

The response shows that the target platform contains user and mgmt debug networks, supports debugging a MicroBlaze™ processor, and also supports JTAG fallback for the Management Debug Network.

Traditionally, a physical JTAG connection is used to perform hardware debug for Xilinx devices with the Vivado hardware manager. The Vitis unified software platforms also makes use of the Xilinx virtual cable (XVC) for hardware debugging on remote accelerator cards. To take advantage of this capability, the Vitis debugger uses the XVC server, an implementation of the XVC protocol that allows the Vivado hardware manager to connect to a local or remote target device for debug, using the standard Xilinx debug cores like the ILA or the VIO IP.

The Vivado hardware manager, from the Vivado Design Suite or Vivado debug feature, can be running on the target instance or it can be running remotely on a different host. The TCP port on which the XVC server is listening must be accessible to the host running Vivado hardware manager. To connect the Vivado hardware manager to XVC server on the target, the following steps should be followed on the machine hosting the Vivado tools:

1. Launch the Vivado debug feature, or the full Vivado Design Suite.
2. Select Open Hardware Manager from the Tasks menu, as shown in the following figure.
   
   
3. Connect to the Vivado tools hw_server, specifying a local or remote connection, and the Host name and Port, as shown below.
   
   
4. Connect to the target instance Virtual JTAG XVC server.
   
   
5. Select the debug_bridge instance from the Hardware window in the Vivado hardware manager.
   Specify the probes file (.ltx) for your design adding it to the Probes > File entry in the Hardware Device Properties window. Adding the probes file refreshes the hardware device, and Hardware window should now show the debug cores in your design.

   TIP: If the kernel has debug cores as specified in Enabling Kernels for Debugging with Chipscope, the probes file (.ltx) is written out during the implementation of the kernel by the Vivado tool.
6. The Vivado hardware manager can now be used to debug the kernels running on the Vitis software platform. Arm the ILA cores in your kernels and run your host application.
   
   

Hardware debug for the Alveo Data Center accelerator cards typically uses the XVC-over-PCIe connection due to the inaccessibility of the physical card, and the JTAG connector on the card. While XVC-over-PCIe allows you to remotely debug your application running on the target platform, certain conditions such as AXI interconnect system hangs can prevent you from accessing the hardware debug functionality that depends on these PCIe/AXI features. Being able to debug these kinds of conditions is especially important for platform designers.

The JTAG Fallback feature is designed to provide access to debug networks that were previously only accessible through XVC-over-PCIe. The JTAG Fallback feature can be enabled without having to change the XVC-over-PCIe-based debug network in the platform design.

On the host side, when the Vivado hardware manager user connects through the hw_server to a JTAG cable that is connected to the physical JTAG pins of the accelerator card, or device under test (DUT), the hw_server disables the XVC-over-PCIe pathway to the hardware. This lets you use the XVC-over-PCIe cable as your primary debug path, but enable debug over the JTAG cable directly when it is required in certain situations. When you disconnect from the JTAG cable, the hw_server re-enables the XVC-over-PCIe pathway to the hardware.

Well-designed kernels and modules report issues through the kernel ring buffer. This is also true for Vitis technology modules that allow you to debug the interaction with the accelerator board on the lowest Linux level.

The dmesg utility is a Linux tool that lets you read the kernel ring buffer. The kernel ring buffer holds kernel information messages in a circular buffer. A circular buffer of fixed size is used to limit the resource requirements by overwriting the oldest entry with the next incoming message.

In the Vitis technology, the xocl module and xclmgmt driver modules write informational messages to the ring buffer. Thus, for an application hang, crash, or any unexpected behavior (like being unable to program the bitstream, etc.), the dmesg tool should be used to check the ring buffer.

The following image shows the layers of the software platform associated with the target platform.

To review messages from the Linux tool, you should first clear the ring buffer:

sudo dmesg -c

This flushes all messages from the ring buffer and makes it easier to spot messages from the xocl and xclmgmt. After that, start your application and run dmesg in another terminal.

sudo dmesg

The dmesg utility prints a record shown in the following example:

In the example shown above, the AXI Firewall 2 has tripped, which is better examined using the xbutil utility.

The Xilinx board utility (xbutil) is a powerful standalone command line utility that can be used to debug lower level hardware/software interaction issues. A full description of this utility can be found in xbutil Utility.

The user errors that typically create application hangs are listed below:

• Read-before-write in 5.0+ target platforms causes a Memory Interface Generator error correction code (MIG ECC) error. This is typically a user error. For example, this error might occur when a kernel is expected to write 4 KB of data in DDR, but it produces only 1 KB of data, and then try to transfer the full 4 KB of data to the host. It can also happen if you supply a 1 KB buffer to a kernel, but the kernel tries to read 4 KB of data.
• An ECC read-before-write error also occurs if no data has been written to a memory location as the last bitstream download which results in MIG initialization, but a read request is made for that same memory location. ECC errors stall the affected MIG because kernels are usually not able to handle this error. This can manifest in two different ways:
  1. The CU might hang or stall because it cannot handle this error while reading or writing to or from the affected MIG. The xbutil query shows that the CU is stuck in a BUSY state and is not making progress.
  2. The AXI Firewall might trip if a PCIe® DMA request is made to the affected MIG, because the DMA engine is unable to complete the request. AXI Firewall trips result in the Linux kernel driver killing all processes which have opened the device node with the SIGBUS signal. The xbutil query shows if an AXI Firewall has indeed tripped and includes a timestamp.
  If the above hang does not occur, the host code might not read back the correct data. This incorrect data is typically 0s and is located in the last part of the data. It is important to review the host code carefully. One common example is compression, where the size of the compressed data is not known up front, and an application might try to migrate more data to the host than was produced by the kernel.

The Vitis compiler is capable of creating very efficient implementations. In some cases, however, implementation issues can occur. One such case is if a write request is emitted before there is enough data available in the process to complete the write transaction. This can cause deadlock conditions when multiple concurrent kernels are affected by this issue and the write request of a kernel depends on the input read being completed.

To avoid these situations, a conservative mode is available on the adapter. In principle, it delays the write request until it has all of the data necessary to complete the write. This mode is enabled during compilation by applying the following --advanced.param option to the v++ compiler:

--advanced.param:compiler.axiDeadLockFree=yes

Because enabling this mode can impact performance, you might prefer to use this as a defensive programming technique where this option is inserted during development and testing and then removed during optimization. You might also want to add this option when the accelerator hangs repeatedly.