diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index 7fb099b06d2..eb9ef8877d8 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/cuda_core/cuda/core/_program.pyx @@ -771,6 +771,8 @@ cdef inline int Program_init(Program self, object code, str code_type, object op assert_type(code, str) if options.extra_sources is not None: raise ValueError("extra_sources is not supported by the NVRTC backend (C++ code_type)") + if options.use_libdevice: + raise ValueError("use_libdevice is not supported by the NVRTC backend (C++ code_type)") # TODO: support pre-loaded headers & include names code_bytes = code.encode() @@ -789,6 +791,8 @@ cdef inline int Program_init(Program self, object code, str code_type, object op assert_type(code, str) if options.extra_sources is not None: raise ValueError("extra_sources is not supported by the PTX backend.") + if options.use_libdevice: + raise ValueError("use_libdevice is not supported by the PTX backend.") code_bytes = code.encode() self._code = code_bytes self._linker = Linker( @@ -835,9 +839,12 @@ cdef inline int Program_init(Program self, object code, str code_type, object op self._linker = None else: + # No use_libdevice check here: this branch is only reached for a + # code_type that is not a backend at all, and an unrecognised + # code_type is the error worth reporting. The per-backend guards + # above mirror the extra_sources ones and are what actually + # enforce "NVVM only", which is what ProgramOptions documents. supported_code_types = tuple(x.value for x in SourceCodeType) - if options.use_libdevice: - raise ValueError("use_libdevice is only supported by the NVVM backend") raise RuntimeError(f"Unsupported {code_type=} ({supported_code_types=})") return 0 diff --git a/cuda_core/docs/source/release/1.2.0-notes.rst b/cuda_core/docs/source/release/1.2.0-notes.rst index 120d2c2a253..bec1df23c21 100644 --- a/cuda_core/docs/source/release/1.2.0-notes.rst +++ b/cuda_core/docs/source/release/1.2.0-notes.rst @@ -73,6 +73,14 @@ Fixes and enhancements Windows, both ``ctypes.CFUNCTYPE`` and ``ctypes.WINFUNCTYPE`` are accepted. (`#2439 `__) +- :class:`Program` now rejects ``use_libdevice=True`` for ``code_type="c++"`` + and ``code_type="ptx"``, as :class:`ProgramOptions` documents. The guard was + written in ``Program_init``'s unrecognised-``code_type`` branch, so the two + real non-NVVM backends accepted the option silently and never linked + libdevice, leaving the caller with undefined-symbol errors at link time. + Conversely, an unrecognised ``code_type`` combined with ``use_libdevice=True`` + now reports the ``code_type``, not libdevice. + Deprecation Notices ------------------- diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 28465425c0e..a5df1dedbde 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -345,6 +345,20 @@ def test_program_init_invalid_code_type(): Program(code, "FORTRAN") +@pytest.mark.agent_authored(model="claude-opus-5") +def test_program_init_invalid_code_type_reports_the_code_type(): + """An unrecognised code_type is the error worth reporting, even when + use_libdevice is set. + + The use_libdevice guard used to live in this branch, so a typo'd + code_type combined with use_libdevice=True reported + "use_libdevice is only supported by the NVVM backend" and never + mentioned that the code_type was the actual problem. + """ + with pytest.raises(RuntimeError, match=r"^Unsupported code_type='fortran'"): + Program("goto 100", "FORTRAN", ProgramOptions(arch="sm_80", use_libdevice=True)) + + def test_program_init_invalid_code_format(): code = 12345 with pytest.raises(TypeError): @@ -718,6 +732,22 @@ def test_cpp_program_with_extra_sources(): Program(code, "c++", options) +@pytest.mark.agent_authored(model="claude-opus-5") +def test_cpp_program_with_use_libdevice(): + """NVRTC has no libdevice loading path. + + The guard used to sit in Program_init's unrecognised-code_type branch, so + "c++" accepted use_libdevice=True silently: Program._use_libdevice stays + False (it is only set inside the nvvm branch), libdevice is never linked, + and the caller finds out from undefined-symbol errors at link time instead + of from the up-front ValueError ProgramOptions documents. + """ + code = 'extern "C" __global__ void my_kernel(){}' + options = ProgramOptions(use_libdevice=True) + with pytest.raises(ValueError, match="use_libdevice is not supported by the NVRTC backend"): + Program(code, "c++", options) + + def test_program_options_as_bytes_nvrtc(): """Test ProgramOptions.as_bytes() for NVRTC backend""" options = ProgramOptions(arch="sm_80", debug=True, lineinfo=True, ftz=True) @@ -831,6 +861,14 @@ def test_ptx_program_extra_sources_unsupported(ptx_code_object): Program(ptx_code_object.code.decode(), "ptx", options) +@pytest.mark.agent_authored(model="claude-opus-5") +def test_ptx_program_use_libdevice_unsupported(ptx_code_object): + """PTX goes through the Linker, which has no libdevice loading path.""" + options = ProgramOptions(use_libdevice=True) + with pytest.raises(ValueError, match="use_libdevice is not supported by the PTX backend"): + Program(ptx_code_object.code.decode(), "ptx", options) + + def test_ptx_program_handle_is_linker_handle(init_cuda, ptx_code_object): """Program.handle for the PTX backend delegates to the linker handle.""" program = Program(ptx_code_object.code.decode(), "ptx")