fix(core): reject use_libdevice on the NVRTC and PTX backends

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

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 <https://github.com/NVIDIA/cuda-python/issues/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
 -------------------
 

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")