[BUG]: Compiling PTX with undefined extern succeeds and produces empty kernel

[gmarkall]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Silent Failure

Component

cuda.core

Describe the bug

Attempting to compile PTX with an undefined extern function seems to succeed and produce a cubin with an empty kernel. I am doing this like:

from cuda.core.experimental import Program, ProgramOptions

options = ProgramOptions(arch="sm_75")
program = Program(code, code_type="ptx", options=options)
cubin = program.compile("cubin")

The same PTX compiled with ptxas results in a failure to compile:

$ ptxas -arch sm_75 add_float16.ptx 
ptxas fatal   : Unresolved extern function '_ZplRK6__halfS1__1'

Am I using the Program.compile() interface incorrectly?

How to Reproduce

Run the attached file:

python cudapy.py

This will produce a cubin with an empty kernel, and show ptxas raises an error for the same PTX.

cudapy.py

Expected behavior

An exception or some other error raised by Program.compile()

Operating System

Ubuntu Linux

nvidia-smi output

No response

[leofang]

This is odd 🤔 Internally, if compiling to PTX is requested, we dispatch to the linker which would invoke either nvJitLink (if exists) or cuLink:

cuda-python/cuda_core/cuda/core/experimental/_program.py

Lines 408 to 413 in 1137e15

	elif code_type == "ptx":
	assert_type(code, str)
	self._linker = Linker(
	ObjectCode._init(code.encode(), code_type), options=self._translate_program_options(options)
	)
	self._backend = self._linker.backend

and I'd assume that either linker would yell because of lack of extern definition. It sounds like the linker decides to just prune undefined externs without raising an error?!

[isVoid]

compile doesn't invoke the linker under the hood. The get_ltoir method from library returns the nvvm compiled result only:

https://github.com/NVIDIA/numba-cuda/blob/main/numba_cuda/numba/cuda/codegen.py#L267-L283

[gmarkall]

@isVoid This is compile from cuda.core.experimental.Program though, not numba.cuda.compile.