Debugger: local variables

[akharche]

Implement info locals for all the kernel arguments and local variables.

Expected:

• Debugger shows all the local variables names inside the numba-dppy kernel when a user calls info locals.
• Debugger shows the values of these local variables when a user calls info locals or print variable.
• Debugger shows a type of the local variables ptype variable.
• User can reset local variable value.

Current state and issues:

1. Some kernel arguments are optimized out when we set optimization level to 1 (by default). Disabling optimization ("-O0") causes segfault due to unsafe address space cast. Solution is in progress: Removing kernel wrapper function #364
2. Local variables created inside a kernel don't have a LLVM debug tag Debugger: DWARF tag is not generated for local kernel variables #320. Temporary workaround was made to Numba (Remove checking on user variable to emit all debug info numba#217). TODO: update results with Numba 0.53.
3. For some reason GDB with Numba shows variable values only on the second entry to the compiled function. Double function call in Numba #382

[akharche]

Numba CPU info locals:

@njit(debug=True)
def foo(arg):
    l1 = arg + 6        	# numba-kernel-breakpoint
    l2 = arg * 5.43             # second-line
    l3 = (arg, l1, l2, "bar")   # third-line
    print(arg, l1, l2, l3)     # fourth-line

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

gdb output for Numba-jit function:

Printing arguments of jit-function is confusing, they are represented as IR lowered values:

[akharche]

Numba "info locals" returns local values in not readable view. It's impossible to understand the value of arguments

[PokhodenkoSA]

Strange behavior:
foo(arg) called twice: with arg = 0 and with arg = 987 (any value passed to foo()). It is unexpected for user. But Numba docs contain example with same behavior.

[akharche]

Forcing Numba to emit debug info for all the variables (pr) gives the following result on GPU:

@dppy.kernel
def data_parallel_sum(user_a, user_b, user_c):
    i = dppy.get_global_id(0)
    user_d = 10
    user_c[i] = user_a[i] + user_b[i] + user_d

def driver(a, b, c, global_size):
    data_parallel_sum[global_size, dppy.DEFAULT_LOCAL_SIZE](a, b, c)

def main():
    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)
    with dpctl.device_context("opencl:gpu") as gpu_queue:
        driver(a, b, c, global_size)

[akharche]

Currently we see all the user variables names while debugging, but they don't have any value.

[akharche]

With Numba 0.53 local variables created inside the kernel are shown:

[PokhodenkoSA]

Fixed in Numba 0.55.