Debugging the compilation pipeline

Consider the following two examples. numba_dpex/examples/debug/simple_sum.py:

import dpctl
import numpy as np

import numba_dpex as dppy


@dppy.kernel(debug=True)
def data_parallel_sum(a, b, c):
    i = dppy.get_global_id(0)
    c[i] = a[i] + b[i]  # Condition breakpoint location


global_size = 10
N = global_size

a = np.array(np.random.random(N), dtype=np.float32)
b = np.array(np.random.random(N), dtype=np.float32)
c = np.ones_like(a)

device = dpctl.SyclDevice("opencl:gpu")
with dpctl.device_context(device):
    data_parallel_sum[global_size, dppy.DEFAULT_LOCAL_SIZE](a, b, c)

print("Done...")

Example of default Numba execution on a CPU:

from numba import njit


@njit(debug=True)
def foo(arg):
    l1 = arg + 6
    l2 = arg * 5.43
    l3 = (arg, l1, l2, "bar")
    print(arg, l1, l2, l3)


def main():
    result = foo(987)
    print(result)


if __name__ == "__main__":
    main()

Getting the DWARF from binary file

In order to get the DWARF information from a specific kernel, just enable IGC_ShaderDumpEnable variable. IGC will write number of dumps into /tmp/IntelIGC.

export IGC_ShaderDumpEnable=1

To read the DWARF of a kernel, a copy of the IGC generated kernel binary is needed. Run the Python script in a GDB debugger mode, and set a breakpoint in the kernel:

$ export NUMBA_OPT=1
$ gdb-oneapi -q --args python simple_sum.py
(gdb) break simple_sum.py:22
(gdb) run

Once the breakpoint is hit, the kernel has been generated and offloaded. At that point, the IGFX driver (i.e. our GDB debugger driver) has copied the kernel into a file, and saved it at /tmp. All files saved at /tmp/IntelIGC/python_xxx/

Dump generated DWARF from the kernel binary (elf) via llvm-dwarfdump tool:

llvm-dwarfdump xxx.elf

Getting the DWARF from Numba assembly (for njit)

Setting Numba environment variable NUMBA_DUMP_ASSEMBLY dumps the native assembly code of compiled functions.

NUMBA_DUMP_ASSEMBLY=1 python njit-basic.py > njit-basic-asm.txt

Clear file from prints and unrecognized characters (-, =, !, ) Get compiled from assembler:

as -o o njit-basic-asm.txt

Get dwarf with objdump:

objdump -W o > o_dwarf

See also:

• NUMBA_DUMP_ASSEMBLY