[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.

[leofang]

We should see an error raised mentioning "undefined reference"

cuda-python/cuda_core/tests/test_linker.py

Lines 147 to 168 in fa07cf5

	def test_linker_get_error_log(compile_ptx_functions):
	options = LinkerOptions(name="ABC", arch=ARCH)
	replacement_kernel = """
	extern __device__ int Z();
	extern __device__ int C(int a, int b);
	__global__ void A() { int result = C(Z(), 1);}
	"""
	dummy_program = Program(
	replacement_kernel, "c++", ProgramOptions(name="CBA", relocatable_device_code=True)
	).compile("ptx")
	linker = Linker(dummy_program, *(compile_ptx_functions[1:]), options=options)
	try:
	linker.link("cubin")
	except (nvJitLinkError, CUDAError):
	log = linker.get_error_log()
	assert isinstance(log, str)
	# TODO when 4902246 is addressed, we can update this to cover nvjitlink as well
	# The error is coming from the input object that's being linked (CBA), not the output object (ABC).
	if is_culink_backend:
	assert log.rstrip("\x00") == "error : Undefined reference to '_Z1Zv' in 'CBA'"

Not sure why it does not happen. Still digging...

[leofang]

@gmarkall I see that if we pass

options = ProgramOptions(arch="sm_75", debug=True)

instead of

options = ProgramOptions(arch="sm_75")

nvJitLink will raise the expected error:

Traceback (most recent call last):
  File "/local/home/leof/dev/debug_920.py", line 95, in <module>
    cubin = program.compile("cubin")
            ^^^^^^^^^^^^^^^^^^^^^^^^
  File "/local/home/leof/dev/cuda-python/cuda_core/cuda/core/experimental/_program.py", line 519, in compile
    return self._linker.link(target_type)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/local/home/leof/dev/cuda-python/cuda_core/cuda/core/experimental/_linker.py", line 448, in link
    _nvjitlink.complete(self._mnff.handle)
  File "cuda/bindings/nvjitlink.pyx", line 156, in cuda.bindings.nvjitlink.complete
  File "cuda/bindings/nvjitlink.pyx", line 166, in cuda.bindings.nvjitlink.complete
  File "cuda/bindings/nvjitlink.pyx", line 66, in cuda.bindings.nvjitlink.check_status
cuda.bindings.nvjitlink.nvJitLinkError: ERROR_INTERNAL (6)
Linker error log: error   : Undefined reference to '_ZplRK6__halfS1__1' in ''
ERROR 9: finish

Do you think this is a bug in nvJitLink?

[gmarkall]

Do you think this is a bug in nvJitLink?

I think so. If I compile and run

#include <iostream>
#include <cstring>
#include "nvJitLink.h"

const char* code = R"PTX(
.version 9.0
.target sm_75
.address_size 64

	// .globl	add
.extern .func  (.param .b32 func_retval0) _ZplRK6__halfS1__1
(
	.param .b64 _ZplRK6__halfS1__1_param_0,
	.param .b64 _ZplRK6__halfS1__1_param_1,
	.param .b64 _ZplRK6__halfS1__1_param_2
)
;
.common .global .align 8 .u64 _ZN08NumbaEnv8__main__3addB2v1B92cw51cXTLSUwv1sDUaKthrqNgqqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE5ArrayIDhLi1E1C7mutable7alignedEDhDh;
.common .global .align 8 .u64 _ZN08NumbaEnv5numba4cuda9_internal9cuda_fp1625_lower__ZplRK6__halfS1__112_3clocals_3e25_ZplRK6__halfS1__1_callerB2v2B92cw51cXTLSUwv1sDUaKthrqNgqqmjgOR3W3CwAkMXLaJtQYkOIgxJU0gCqOkEJoHkbttqdVhoqlspQGNFHSgJ5BnXagIAE10float16_2a10float16_2a;

.visible .func  (.param .b64 func_retval0) add(
	.param .b64 add_param_0,
	.param .b64 add_param_1,
	.param .b64 add_param_2,
	.param .b64 add_param_3,
	.param .b64 add_param_4,
	.param .b64 add_param_5,
	.param .b64 add_param_6,
	.param .b32 add_param_7,
	.param .b32 add_param_8
)
{
	.local .align 2 .b8 	__local_depot0[6];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<3>;
	.reg .b16 	%rs<5>;
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<9>;


	mov.u64 	%SPL, __local_depot0;
	cvta.local.u64 	%SP, %SPL;
	ld.param.u64 	%rd2, [add_param_4];
	ld.param.u16 	%rs1, [add_param_7];
	ld.param.u16 	%rs2, [add_param_8];
	add.u64 	%rd3, %SP, 0;
	add.u64 	%rd1, %SPL, 0;
	add.u64 	%rd4, %SP, 2;
	add.u64 	%rd5, %SPL, 2;
	add.u64 	%rd6, %SP, 4;
	add.u64 	%rd7, %SPL, 4;
	st.local.u16 	[%rd5], %rs1;
	st.local.u16 	[%rd7], %rs2;
	mov.u16 	%rs3, 0;
	st.local.u16 	[%rd1], %rs3;
	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b64 param0;
	st.param.b64 	[param0+0], %rd3;
	.param .b64 param1;
	st.param.b64 	[param1+0], %rd4;
	.param .b64 param2;
	st.param.b64 	[param2+0], %rd6;
	.param .b32 retval0;
	call.uni (retval0), 
	_ZplRK6__halfS1__1, 
	(
	param0, 
	param1, 
	param2
	);
	ld.param.b32 	%r1, [retval0+0];
	} // callseq 0
	setp.eq.s32 	%p1, %r1, 0;
	@%p1 bra 	$L__BB0_2;

	setp.ne.s32 	%p2, %r1, -2;
	@%p2 bra 	$L__BB0_3;

$L__BB0_2:
	ld.local.u16 	%rs4, [%rd1];
	st.u16 	[%rd2], %rs4;

$L__BB0_3:
	mov.u64 	%rd8, 0;
	st.param.b64 	[func_retval0+0], %rd8;
	ret;
}
)PTX";

#define NVJITLINK_SAFE_CALL(h,x)                                  \
  do {                                                            \
    nvJitLinkResult result = x;                                   \
    if (result != NVJITLINK_SUCCESS) {                            \
      std::cerr << "\nerror: " #x " failed with error "           \
                << result << '\n';                                \
      size_t lsize;                                               \
      result = nvJitLinkGetErrorLogSize(h, &lsize);               \
      if (result == NVJITLINK_SUCCESS && lsize > 0) {             \
        char *log = (char*)malloc(lsize);                         \
    result = nvJitLinkGetErrorLog(h, log);                        \
    if (result == NVJITLINK_SUCCESS) {                            \
      std::cerr << "error: " << log << '\n';                      \
      free(log);                                                  \
    }                                                             \
      }                                                           \
      exit(1);                                                    \
    }                                                             \
  } while(0)

int main(int argc, char **argv) {
  nvJitLinkHandle handle;

  const char *opts[] = {"-arch=sm_75"};
  NVJITLINK_SAFE_CALL(handle, nvJitLinkCreate(&handle, 1, opts));
  NVJITLINK_SAFE_CALL(handle, nvJitLinkAddData(handle, NVJITLINK_INPUT_PTX, (void*)code, strlen(code), "code.ptx"));
  NVJITLINK_SAFE_CALL(handle, nvJitLinkComplete(handle));

  return 0;
}

with

export NVJITLINK_PATH=$CONDA_PREFIX/targets/x86_64-linux
g++ -o cuda-python-920 -I$NVJITLINK_PATH/include -L$NVJITLINK_PATH/lib -lnvJitLink -Wl,-rpath,$NVJITLINK_PATH/lib cuda-python-920.cpp
./cuda-python-920

then I get

$ ./cuda-python-920 

$ echo $?
0

but if I change the options to

  const char *opts[] = {"-arch=sm_75", "-g"};
  NVJITLINK_SAFE_CALL(handle, nvJitLinkCreate(&handle, 2, opts));

then I get

$ ./cuda-python-920 

error: nvJitLinkComplete(handle) failed with error 6
error: error   : Undefined reference to '_ZplRK6__halfS1__1' in 'code.ptx'
ERROR 9: finish

[leofang]

@gmarkall I am scratching my head again because the following script is equivalent to what we test as mentioned above, and it does raise an error:

Traceback (most recent call last):
  File "/local/home/leof/dev/cuda-python/cuda_core/debug_nvjitlink.py", line 64, in <module>
    linker.link(target_type="cubin")
  File "/local/home/leof/dev/cuda-python/cuda_core/cuda/core/experimental/_linker.py", line 483, in link
    _nvjitlink.complete(self._mnff.handle)
  File "cuda/bindings/nvjitlink.pyx", line 165, in cuda.bindings.nvjitlink.complete
  File "cuda/bindings/nvjitlink.pyx", line 175, in cuda.bindings.nvjitlink.complete
  File "cuda/bindings/nvjitlink.pyx", line 75, in cuda.bindings.nvjitlink.check_status
cuda.bindings.nvjitlink.nvJitLinkError: ERROR_INTERNAL (6)
Linker error log: error   : Undefined reference to '_Z1Cii' in ''
error   : Undefined reference to '_Z1Zv' in ''
ERROR 9: finish

but it does not need -g (debug=True) 🤔

from cuda.core.experimental import Device, Linker, LinkerOptions, ObjectCode


ptx_code = b"""
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-36424714
// Cuda compilation tools, release 13.0, V13.0.88
// Based on NVVM 7.0.1
//

.version 9.0
.target sm_75
.address_size 64

	// .globl	_Z1Av
.extern .func  (.param .b32 func_retval0) _Z1Cii
(
	.param .b32 _Z1Cii_param_0,
	.param .b32 _Z1Cii_param_1
)
;
.extern .func  (.param .b32 func_retval0) _Z1Zv
()
;

.visible .entry _Z1Av()
{
	.reg .b32 	%r<3>;


	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b32 retval0;
	call.uni (retval0), 
	_Z1Zv, 
	(
	);
	ld.param.b32 	%r1, [retval0+0];
	} // callseq 0
	{ // callseq 1, 0
	.reg .b32 temp_param_reg;
	.param .b32 param0;
	st.param.b32 	[param0+0], %r1;
	.param .b32 param1;
	st.param.b32 	[param1+0], 1;
	.param .b32 retval0;
	call.uni (retval0), 
	_Z1Cii, 
	(
	param0, 
	param1
	);
	ld.param.b32 	%r2, [retval0+0];
	} // callseq 1
	ret;

}
"""

ptx_obj = ObjectCode.from_ptx(module=ptx_code)
linker = Linker(ptx_obj, options=LinkerOptions(arch="sm_75"))
linker.link(target_type="cubin")

[gmarkall]

The example in the previous comment that doesn't require -g has a kernel in it calling the extern device function, but the original example has a device function calling an extern device function.

So, it looks like this is a difference in the configurations for the link between ptxas and nvJitLink - ptxas seems to try to produce a relocatable cubin by default, and nvJitLink seems to try to produce a fully linked one - So it's probably just discarding the caller in the original example and therefore an error is not being produced. When -g is applied, that probably disables the optimization that prunes dead functions, so the error gets emitted.

[gmarkall]

Adding options like -Xptxas=-c" and -Xptxas=-ewp don't seem to help though. Perhaps nvJitLink is running ptxas in compile-only (-c) mode to create a relocatable cubin, but then also linking the output causing the device function with a call to an undefined function to be dropped anyway.