[cuda] lineinfo for device function

[sklam]

Related gitter discussion starting at here and here.

Reproducer:

# filename: debugme.py
import math
import numpy as np
import numba
from numba import cuda


@cuda.jit(device=True, debug=True)
def expensive(x):
    return math.pow(x, 1.001)


@cuda.jit(device=True, debug=True)
def cheap(x):
    return math.sqrt(x)


@cuda.jit(device=True, debug=True)
def calc(x):
    s1 = expensive(x)
    q = cheap(s1)
    return q + 1


@cuda.jit(device=True, debug=True)
def threadid():
    return cuda.blockDim.x * cuda.blockIdx.x + cuda.threadIdx.x


@cuda.jit(argtypes=[numba.int32[:]], debug=True)
def kernel(arr):
    for i in range(1000):
        where = threadid()
        value = arr[where]
        arr[where] = calc(value)


def main():
    print("SUP")
    z = np.array(range(0, 5000))
    print(z)
    kernel[5000 // 128 + 1, 128](z)
    print(z)

    print(kernel.inspect_llvm())

if __name__ == '__main__':
    main()

Run the above with:

$ nvprof -a instruction_execution --export-profile profile_output python debugme.py

On my OSX 10.13 with CUDA 9.2, we can see some line-info in the profile_output using NVVP:

The "inline-closure" version (this style: https://gist.github.com/c200chromebook/d131322c918b142f14f40c9b3583cc6d) looks like:

@c200chromebook, can you tell us what you see?

[c200chromebook]

Win 2016 Datacenter, CUDA 10.1, numba 0.48.0, numpy 1.18.1, llvmlite 0.31.0 - you don't even get near the profiling.

[c200chromebook]

Using the goofy inline-closure format:

[sklam]

hmm... i think this might be a problem in libNVVM in CUDA 10.1. @gmarkall, do you know if there are known problems regarding debug-symbols in newer CUDA.

[c200chromebook]

Yeah, you'll excuse me, I did very minorly goof - 10.2. Forgot we had upgraded.

[gmarkall]

@sklam If you're referring to the appearance of the message "Debugging support cannot be enabled when the number of debug compile units is more than 1", it's a message I've bumped into before but hadn't yet got around to looking into... I didn't realise it ever did work with 9.2 though.

[stuartarchibald]

I've also seen that message a fair few times.

[sklam]

@stuartarchibald, on what hardwares and CUDA versions?

[gmarkall]

I've seen it with CUDA 10.0 and a V100.

[stuartarchibald]

@stuartarchibald, on what hardwares and CUDA versions?

Stretching memory, but CUDA 10.2 and a GTX970.

[gmarkall]

hmm... i think this might be a problem in libNVVM in CUDA 10.1. @gmarkall, do you know if there are known problems regarding debug-symbols in newer CUDA.

I just noticed that it's documented that only a single module is supported: https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#source-level-debugging-support:

Source level debugging is supported only for a single debug compile unit. If there are multiple input NVVM IR modules, at most one module may have a single debug compile unit.

[sklam]

Hm... we can possibly prelink the module in llvm before sending it over to nvvm.

[c200chromebook]

Is this one a fundamental limitation or should the prelink workaround fix?

[sklam]

@gmarkall, the doc change first happened in 10.0 (https://docs.nvidia.com/cuda/archive/10.0/nvvm-ir-spec/index.html#source-level-debugging-support). In 9.2 (https://docs.nvidia.com/cuda/archive/9.2/nvvm-ir-spec/index.html#source-level-debugging-support), there were no restriction on "a single debug compile unit" and it works.

Why is this restriction needed? I would think that one can use LLVM to link the modules together and the debug symbol will merge fine. If it is not that simple, I don't know what Numba can do to address this problem.

@c200chromebook is reporting (at https://gitter.im/numba/numba?at=5f35a6314bb58d3181951aea) that this affects a single kernel case in Numba. Which is odd because there should only be one LLVM module. Unless NVVM is counting the libdevice that we are linking in.

[c200chromebook]

See also the below, which has two items to compile, thus causing a fail.

from numba import cuda
import numba as nb

@cuda.jit(nb.void(nb.int32[:], nb.int32[:]), debug=True)
def whatever(inp, outp):
    outp[0] = 1 if inp[0] in (2, 3) else 3


whatever[1, 1](cuda.device_array(1), cuda.device_array(1))

yields

Error: Debugging support cannot be enabled when number of debug compile units is more than 1
NVVM_ERROR_COMPILATION

This is an idiosyncrasy of the tuple, as you don't get it with:

@cuda.jit(nb.void(nb.int32[:], nb.int32[:]), debug=True)
def whatever(inp, outp):
    outp[0] = 1 if inp[0] == 2 or inp[0] ==3 else 3

[gmarkall]

On master with CUDA 11.2 (NVVM 7.0) line info seems to be available again with debug=True for the reproducer in the issue description:

[gmarkall]

See also the below, which has two items to compile, thus causing a fail.

from numba import cuda
import numba as nb

@cuda.jit(nb.void(nb.int32[:], nb.int32[:]), debug=True)
def whatever(inp, outp):
    outp[0] = 1 if inp[0] in (2, 3) else 3


whatever[1, 1](cuda.device_array(1), cuda.device_array(1))

yields

Error: Debugging support cannot be enabled when number of debug compile units is more than 1
NVVM_ERROR_COMPILATION

This reproduces for me with CUDA 10.1, but the issue appears to be resolved with CUDA 11.2.

[gmarkall]

To summarise where I think we are:

• Our debuginfo generation is better for NVVM 7.0 than it is for NVVM 3.4, due to using the exact same mechanism as the CPU target.
• For a presently-unknown reason, the "Error: Debugging support cannot be enabled when number of debug compile units is more than 1" is not triggered with NVVM 7.0, though this is still unsupported: https://docs.nvidia.com/cuda/archive/11.2.1/nvvm-ir-spec/index.html#source-level-debugging-support
• See also the PR to add lineinfo generation only: Add lineinfo flag to PTX and SASS compilation #6802

@sklam asked earlier:

Why is this restriction needed? I would think that one can use LLVM to link the modules together and the debug symbol will merge fine. If it is not that simple, I don't know what Numba can do to address this problem.

I don't know the reason for the restriction. I believe that we can conform with this restriction in Numba by changing

numba/numba/cuda/codegen.py

Lines 114 to 115 in 81f7082

	irs = [str(mod) for mod in self.modules]
	ptx = nvvm.llvm_to_ptx(irs, **options)

so that instead of passing a list of the IR for all modules into llvm_to_ptx, we call it once for each module's IR, then link the resulting PTX for each module one-by-one when linking the cubin. So here:

numba/numba/cuda/codegen.py

Line 139 in 81f7082

linker.add_ptx(ptx.encode())

we'd have multiple PTXs to add.

I hadn't done this by default so far because it seems to result in less inlining and less efficient code in the non-debug case (I did try it, but couldn't get the linker to inline the PTX as eagerly as NVVM seems to inline IR for device functions). For debug, this isn't a problem though, as we're not going to be expecting much / any inlining.

We have this hack in the code:

numba/numba/cuda/cudadrv/nvvm.py

Lines 708 to 728 in 81f7082

	# XXX remove debug_pubnames seems to be necessary sometimes
	return patch_ptx_debug_pubnames(ptx)
	def patch_ptx_debug_pubnames(ptx):
	"""
	Patch PTX to workaround .debug_pubnames NVVM error::
	ptxas fatal : Internal error: overlapping non-identical data
	"""
	while True:
	# Repeatedly remove debug_pubnames sections
	start = ptx.find(b'.section .debug_pubnames')
	if start < 0:
	break
	stop = ptx.find(b'}', start)
	if stop < 0:
	raise ValueError('missing "}"')
	ptx = ptx[:start] + ptx[stop + 1:]
	return ptx

which is intended to workaround the fact we are passing in multiple modules when debug=True. If the above changes to pass one module to NVVM at a time are implemented, I believe that this workaround could be removed (and maybe debug info in CUDA would work a bit better than it presently does).

[gmarkall]

instead of passing a list of the IR for all modules into llvm_to_ptx, we call it once for each module's IR, then link the resulting PTX for each module one-by-one when linking the cubin.

I've tried implementing this strategy and removing the debug_pubnames patching in: gmarkall@ce840bc

It fixes the failure to compile with NVVM 3.4. Things seem to work as well as they did before with NVVM 7.0. Stepping line-by-line seems like it might be more stable, but I'm not sure if that's just my perception - I need to do some more careful experiments.

[gmarkall]

PR #6841 includes a fix for the message "Error: Debugging support cannot be enabled when number of debug compile units is more than 1" using the strategy described above - since lineinfo is working with NVVM 7.0, and the error message is resolved with NVVM 3.4, I will consider this issue resolved once that PR is merged.