fix(core): accept the documented str symbol_mapping in ObjectCode.get_kernel

cuda_core/cuda/core/_module.pyx

@@ -636,7 +636,7 @@ cdef class ObjectCode:
         )
 
     @classmethod
-    def _init(cls, module, code_type, *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def _init(cls, module, code_type, *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         assert code_type in _supported_code_type, f"{code_type=} is not supported"
         cdef ObjectCode self = ObjectCode.__new__(ObjectCode)
 
@@ -651,7 +651,13 @@ cdef class ObjectCode:
             self._module = fspath(module)
         else:
             self._module = module
-        self._sym_map = {} if symbol_mapping is None else symbol_mapping
+        # Normalise here so `_sym_map` is homogeneous: `Program.compile` supplies
+        # bytes (nvrtcGetLoweredName returns const char*), while a hand-built
+        # mapping is documented as str. Storing one type removes the need to
+        # probe the value's type on every `get_kernel` call.
+        self._sym_map = {} if symbol_mapping is None else {
+            k: (v.encode() if isinstance(v, str) else v) for k, v in symbol_mapping.items()
+        }
         self._name = name if name else ""
 
         return self
@@ -664,7 +670,7 @@ cdef class ObjectCode:
         return ObjectCode._reduce_helper, (self._module, self._code_type, self._name, self._sym_map)
 
     @staticmethod
-    def from_cubin(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def from_cubin(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         """Create an :class:`ObjectCode` instance from an existing cubin.
 
         Parameters
@@ -683,7 +689,7 @@ cdef class ObjectCode:
         return ObjectCode._init(module, ObjectCodeFormatType.CUBIN, name=name, symbol_mapping=symbol_mapping)
 
     @staticmethod
-    def from_ptx(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def from_ptx(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         """Create an :class:`ObjectCode` instance from an existing PTX.
 
         Parameters
@@ -702,7 +708,7 @@ cdef class ObjectCode:
         return ObjectCode._init(module, ObjectCodeFormatType.PTX, name=name, symbol_mapping=symbol_mapping)
 
     @staticmethod
-    def from_ltoir(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def from_ltoir(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         """Create an :class:`ObjectCode` instance from an existing LTOIR.
 
         Parameters
@@ -721,7 +727,7 @@ cdef class ObjectCode:
         return ObjectCode._init(module, ObjectCodeFormatType.LTOIR, name=name, symbol_mapping=symbol_mapping)
 
     @staticmethod
-    def from_fatbin(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def from_fatbin(module: bytes | str | PathLike[str], *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         """Create an :class:`ObjectCode` instance from an existing fatbin.
 
         Parameters
@@ -740,7 +746,7 @@ cdef class ObjectCode:
         return ObjectCode._init(module, ObjectCodeFormatType.FATBIN, name=name, symbol_mapping=symbol_mapping)
 
     @staticmethod
-    def from_object(module: bytes | str, *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def from_object(module: bytes | str, *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         """Create an :class:`ObjectCode` instance from an existing object code.
 
         Parameters
@@ -758,7 +764,7 @@ cdef class ObjectCode:
         return ObjectCode._init(module, ObjectCodeFormatType.OBJECT, name=name, symbol_mapping=symbol_mapping)
 
     @staticmethod
-    def from_library(module: bytes | str, *, name: str = "", symbol_mapping: dict[str, str] | None = None) -> ObjectCode:
+    def from_library(module: bytes | str, *, name: str = "", symbol_mapping: dict[str, str | bytes] | None = None) -> ObjectCode:
         """Create an :class:`ObjectCode` instance from an existing library.
 
         Parameters
@@ -802,6 +808,8 @@ cdef class ObjectCode:
         try:
             name = self._sym_map[name]
         except KeyError:
+            # Not a mangled-name lookup, so `name` is still whatever the caller
+            # passed and may need encoding before it reaches `<const char*>`.
             if isinstance(name, str):
                 name = name.encode()
 
@@ -845,8 +853,12 @@ cdef class ObjectCode:
         return self._code_type
 
     @property
-    def symbol_mapping(self) -> dict[str, str]:
-        """Return a copy of the symbol mapping dictionary."""
+    def symbol_mapping(self) -> dict[str, bytes]:
+        """Return a copy of the symbol mapping dictionary.
+
+        Values are always ``bytes``: a ``str`` passed to a factory method is
+        encoded when it is stored.
+        """
         return dict(self._sym_map)
 
     @property

cuda_core/cuda/core/_program.pyx

@@ -903,7 +903,12 @@ cdef object _nvrtc_compile_and_extract(
             name_bytes = n.encode() if isinstance(n, str) else n
             name_ptr = <const char*>name_bytes
             HANDLE_RETURN_NVRTC(prog, cynvrtc.nvrtcGetLoweredName(prog, name_ptr, &lowered_name))
-            symbol_mapping[n] = lowered_name if lowered_name != NULL else None
+            if lowered_name == NULL:
+                # HANDLE_RETURN_NVRTC above already raises on failure, so this is
+                # defensive. Storing None would reach <const char*>None in
+                # ObjectCode.get_kernel and segfault, so refuse it here instead.
+                raise RuntimeError(f"nvrtcGetLoweredName returned no lowered name for {n!r}")
+            symbol_mapping[n] = lowered_name
 
     # Get compilation log if requested
     if logs is not None:

cuda_core/docs/source/release/1.2.0-notes.rst

@@ -73,6 +73,13 @@ Fixes and enhancements
   Windows, both ``ctypes.CFUNCTYPE`` and ``ctypes.WINFUNCTYPE`` are accepted.
   (`#2439 <https://github.com/NVIDIA/cuda-python/issues/2439>`__)
 
+- :meth:`ObjectCode.get_kernel` now accepts a ``symbol_mapping`` whose values
+  are ``str``, as :meth:`ObjectCode.from_cubin` and its siblings document
+  (``dict[str, str]``). The mapped name was only encoded on the *miss* path, so
+  a hand-built mapping raised ``TypeError: expected bytes, str found`` for
+  exactly the names it was supposed to translate. Mappings produced by
+  :meth:`Program.compile` are unaffected -- their values are already ``bytes``.
+
 Deprecation Notices
 -------------------
 

cuda_core/tests/test_module.py

@@ -369,6 +369,27 @@ def test_object_code_load_cubin(get_saxpy_kernel_cubin):
     mod.get_kernel("saxpy<double>")  # force loading
 
 
+@pytest.mark.agent_authored(model="claude-opus-5")
+def test_object_code_symbol_mapping_accepts_str_values(get_saxpy_kernel_cubin):
+    """``symbol_mapping`` is annotated and documented ``dict[str, str]``.
+
+    ``Program.compile`` stores the lowered names as ``bytes`` (they come from
+    ``nvrtcGetLoweredName``), so every existing test round-trips
+    ``mod.symbol_mapping`` unchanged and the documented ``str`` form is never
+    exercised. On a mapping *hit* the value went straight to ``<const char*>``
+    without being encoded, so a hand-built ``dict[str, str]`` raised
+    ``TypeError: expected bytes, str found`` -- the mapping worked only for
+    names it did not map.
+    """
+    _, mod = get_saxpy_kernel_cubin
+    cubin = mod.code
+    str_sym_map = {k: v.decode() if isinstance(v, bytes) else v for k, v in mod.symbol_mapping.items()}
+    assert all(isinstance(v, str) for v in str_sym_map.values())
+
+    obj = ObjectCode.from_cubin(cubin, symbol_mapping=str_sym_map)
+    obj.get_kernel("saxpy<double>")  # force loading through the mapped name
+
+
 def test_object_code_load_cubin_from_file(get_saxpy_kernel_cubin, tmp_path, convert_path):
     _, mod = get_saxpy_kernel_cubin
     cubin = mod.code