Merge branch 'main' into error_explanations

Author: rwgk

cuda_bindings/cuda/bindings/_lib/cyruntime/cyruntime.pxd.in

@@ -108,9 +108,9 @@ from libcpp cimport bool
 {{if 'cudaCreateSurfaceObject' in found_functions}}cdef cudaError_t _cudaCreateSurfaceObject(cudaSurfaceObject_t* pSurfObject, const cudaResourceDesc* pResDesc) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
 {{if 'cudaGetTextureObjectResourceDesc' in found_functions}}cdef cudaError_t _cudaGetTextureObjectResourceDesc(cudaResourceDesc* pResDesc, cudaTextureObject_t texObject) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
 {{if 'cudaGraphicsEGLRegisterImage' in found_functions}}cdef cudaError_t _cudaGraphicsEGLRegisterImage(cudaGraphicsResource_t* pCudaResource, EGLImageKHR image, unsigned int flags) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
-{{if 'cudaEGLStreamProducerPresentFrame' in found_functions}}cdef cudaError_t _cudaEGLStreamProducerPresentFrame(cudaEglStreamConnection* conn, cudaEglFrame eglframe, cudaStream_t* pStream) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
-{{if 'cudaEGLStreamProducerReturnFrame' in found_functions}}cdef cudaError_t _cudaEGLStreamProducerReturnFrame(cudaEglStreamConnection* conn, cudaEglFrame* eglframe, cudaStream_t* pStream) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
-{{if 'cudaGraphicsResourceGetMappedEglFrame' in found_functions}}cdef cudaError_t _cudaGraphicsResourceGetMappedEglFrame(cudaEglFrame* eglFrame, cudaGraphicsResource_t resource, unsigned int index, unsigned int mipLevel) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
+{{if True}}cdef cudaError_t _cudaEGLStreamProducerPresentFrame(cudaEglStreamConnection* conn, cudaEglFrame eglframe, cudaStream_t* pStream) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
+{{if True}}cdef cudaError_t _cudaEGLStreamProducerReturnFrame(cudaEglStreamConnection* conn, cudaEglFrame* eglframe, cudaStream_t* pStream) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
+{{if True}}cdef cudaError_t _cudaGraphicsResourceGetMappedEglFrame(cudaEglFrame* eglFrame, cudaGraphicsResource_t resource, unsigned int index, unsigned int mipLevel) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
 {{if True}}cdef cudaError_t _cudaVDPAUSetVDPAUDevice(int device, VdpDevice vdpDevice, VdpGetProcAddress* vdpGetProcAddress) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
 {{if 'cudaArrayGetMemoryRequirements' in found_functions}}cdef cudaError_t _cudaArrayGetMemoryRequirements(cudaArrayMemoryRequirements* memoryRequirements, cudaArray_t array, int device) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}
 {{if 'cudaMipmappedArrayGetMemoryRequirements' in found_functions}}cdef cudaError_t _cudaMipmappedArrayGetMemoryRequirements(cudaArrayMemoryRequirements* memoryRequirements, cudaMipmappedArray_t mipmap, int device) except ?cudaErrorCallRequiresNewerDriver nogil{{endif}}

cuda_bindings/cuda/bindings/_lib/cyruntime/cyruntime.pyx.in

@@ -2206,6 +2206,7 @@ cdef cudaError_t _cudaGetTextureObjectResourceDesc(cudaResourceDesc* pResDesc, c
     return err
 
 {{endif}}
+{{if True}}
 
 cdef cudaError_t _cudaEGLStreamProducerPresentFrame(cudaEglStreamConnection* conn, cudaEglFrame eglframe, cudaStream_t* pStream) except ?cudaErrorCallRequiresNewerDriver nogil:
     cdef cudaError_t err = cudaSuccess
@@ -2222,6 +2223,9 @@ cdef cudaError_t _cudaEGLStreamProducerPresentFrame(cudaEglStreamConnection* con
         _setLastError(err)
     return err
 
+{{endif}}
+{{if True}}
+
 cdef cudaError_t _cudaEGLStreamProducerReturnFrame(cudaEglStreamConnection* conn, cudaEglFrame* eglframe, cudaStream_t* pStream) except ?cudaErrorCallRequiresNewerDriver nogil:
     cdef cudaError_t err = cudaSuccess
     err = m_global.lazyInitContextState()
@@ -2242,6 +2246,9 @@ cdef cudaError_t _cudaEGLStreamProducerReturnFrame(cudaEglStreamConnection* conn
         return err
     return err
 
+{{endif}}
+{{if True}}
+
 cdef cudaError_t _cudaGraphicsResourceGetMappedEglFrame(cudaEglFrame* eglFrame, cudaGraphicsResource_t resource, unsigned int index, unsigned int mipLevel) except ?cudaErrorCallRequiresNewerDriver nogil:
     cdef cudaError_t err = cudaSuccess
     err = m_global.lazyInitContextState()
@@ -2259,9 +2266,13 @@ cdef cudaError_t _cudaGraphicsResourceGetMappedEglFrame(cudaEglFrame* eglFrame,
         return err
     return err
 
+{{endif}}
+{{if True}}
+
 cdef cudaError_t _cudaVDPAUSetVDPAUDevice(int device, VdpDevice vdpDevice, VdpGetProcAddress* vdpGetProcAddress) except ?cudaErrorCallRequiresNewerDriver nogil:
     return cudaErrorNotSupported
 
+{{endif}}
 {{if 'cudaArrayGetMemoryRequirements' in found_functions}}
 
 cdef cudaError_t _cudaArrayGetMemoryRequirements(cudaArrayMemoryRequirements* memoryRequirements, cudaArray_t array, int device) except ?cudaErrorCallRequiresNewerDriver nogil:

cuda_bindings/docs/source/release/12.X.Y-notes.rst

@@ -9,3 +9,4 @@ Highlights
 
 * The ``cuda.bindings.nvvm`` Python module was added, wrapping the
   `libNVVM C API <https://docs.nvidia.com/cuda/libnvvm-api/>`_.
+* Source build error checking added for missing required headers

cuda_bindings/setup.py

@@ -48,8 +48,11 @@
 # ----------------------------------------------------------------------
 # Parse user-provided CUDA headers
 
-header_dict = {
-    "driver": ["cuda.h", "cudaProfiler.h", "cudaEGL.h", "cudaGL.h", "cudaVDPAU.h"],
+required_headers = {
+    "driver": [
+        "cuda.h",
+        "cudaProfiler.h",
+    ],
     "runtime": [
         "driver_types.h",
         "vector_types.h",
@@ -61,32 +64,42 @@
         "device_types.h",
         "driver_functions.h",
         "cuda_profiler_api.h",
-        "cuda_egl_interop.h",
-        "cuda_gl_interop.h",
-        "cuda_vdpau_interop.h",
     ],
-    "nvrtc": ["nvrtc.h"],
+    "nvrtc": [
+        "nvrtc.h",
+    ],
+    # During compilation, Cython will reference C headers that are not
+    # explicitly parsed above. These are the known dependencies:
+    #
+    # - crt/host_defines.h
+    # - builtin_types.h
+    # - cuda_device_runtime_api.h
 }
 
-replace = {
-    " __device_builtin__ ": " ",
-    "CUDARTAPI ": " ",
-    "typedef __device_builtin__ enum cudaError cudaError_t;": "typedef cudaError cudaError_t;",
-    "typedef __device_builtin__ enum cudaOutputMode cudaOutputMode_t;": "typedef cudaOutputMode cudaOutputMode_t;",
-    "typedef enum cudaError cudaError_t;": "typedef cudaError cudaError_t;",
-    "typedef enum cudaOutputMode cudaOutputMode_t;": "typedef cudaOutputMode cudaOutputMode_t;",
-    "typedef enum cudaDataType_t cudaDataType_t;": "",
-    "typedef enum libraryPropertyType_t libraryPropertyType_t;": "",
-    "  enum ": "   ",
-    ", enum ": ", ",
-    "\\(enum ": "(",
-}
 
-found_types = []
-found_functions = []
-found_values = []
-found_struct = []
-struct_list = {}
+def fetch_header_paths(required_headers, include_path_list):
+    header_dict = {}
+    missing_headers = []
+    for library, header_list in required_headers.items():
+        header_paths = []
+        for header in header_list:
+            path_candidate = [os.path.join(path, header) for path in include_path_list]
+            for path in path_candidate:
+                if os.path.exists(path):
+                    header_paths += [path]
+                    break
+            else:
+                missing_headers += [header]
+
+        # Update dictionary with validated paths to headers
+        header_dict[library] = header_paths
+
+    if missing_headers:
+        error_message = "Couldn't find required headers: "
+        error_message += ", ".join([header for header in missing_headers])
+        raise RuntimeError(f'{error_message}\nIs CUDA_HOME setup correctly? (CUDA_HOME="{CUDA_HOME}")')
+
+    return header_dict
 
 
 class Struct:
@@ -117,52 +130,66 @@ def __repr__(self):
         return f"{self._name}: {self._member_names} with types {self._member_types}"
 
 
-include_path_list = [os.path.join(path, "include") for path in CUDA_HOME]
-print(f'Parsing headers in "{include_path_list}" (Caching = {PARSER_CACHING})')
-for library, header_list in header_dict.items():
-    header_paths = []
-    for header in header_list:
-        path_candidate = [os.path.join(path, header) for path in include_path_list]
-        for path in path_candidate:
-            if os.path.exists(path):
-                header_paths += [path]
-                break
-        if not os.path.exists(path):
-            print(f"Missing header {header}")
-
-    print(f"Parsing {library} headers")
-    parser = CParser(
-        header_paths, cache="./cache_{}".format(library.split(".")[0]) if PARSER_CACHING else None, replace=replace
-    )
+def parse_headers(header_dict):
+    found_types = []
+    found_functions = []
+    found_values = []
+    found_struct = []
+    struct_list = {}
+
+    replace = {
+        " __device_builtin__ ": " ",
+        "CUDARTAPI ": " ",
+        "typedef __device_builtin__ enum cudaError cudaError_t;": "typedef cudaError cudaError_t;",
+        "typedef __device_builtin__ enum cudaOutputMode cudaOutputMode_t;": "typedef cudaOutputMode cudaOutputMode_t;",
+        "typedef enum cudaError cudaError_t;": "typedef cudaError cudaError_t;",
+        "typedef enum cudaOutputMode cudaOutputMode_t;": "typedef cudaOutputMode cudaOutputMode_t;",
+        "typedef enum cudaDataType_t cudaDataType_t;": "",
+        "typedef enum libraryPropertyType_t libraryPropertyType_t;": "",
+        "  enum ": "   ",
+        ", enum ": ", ",
+        "\\(enum ": "(",
+    }
+
+    print(f'Parsing headers in "{include_path_list}" (Caching = {PARSER_CACHING})')
+    for library, header_paths in header_dict.items():
+        print(f"Parsing {library} headers")
+        parser = CParser(
+            header_paths, cache="./cache_{}".format(library.split(".")[0]) if PARSER_CACHING else None, replace=replace
+        )
+
+        if library == "driver":
+            CUDA_VERSION = parser.defs["macros"].get("CUDA_VERSION", "Unknown")
+            print(f"Found CUDA_VERSION: {CUDA_VERSION}")
+
+        # Combine types with others since they sometimes get tangled
+        found_types += {key for key in parser.defs["types"]}
+        found_types += {key for key in parser.defs["structs"]}
+        found_types += {key for key in parser.defs["unions"]}
+        found_types += {key for key in parser.defs["enums"]}
+        found_functions += {key for key in parser.defs["functions"]}
+        found_values += {key for key in parser.defs["values"]}
+
+        for key, value in parser.defs["structs"].items():
+            struct_list[key] = Struct(key, value["members"])
+        for key, value in parser.defs["unions"].items():
+            struct_list[key] = Struct(key, value["members"])
+
+        for key, value in struct_list.items():
+            if key.startswith("anon_union") or key.startswith("anon_struct"):
+                continue
 
-    if library == "driver":
-        CUDA_VERSION = parser.defs["macros"].get("CUDA_VERSION", "Unknown")
-        print(f"Found CUDA_VERSION: {CUDA_VERSION}")
-
-    # Combine types with others since they sometimes get tangled
-    found_types += {key for key in parser.defs["types"]}
-    found_types += {key for key in parser.defs["structs"]}
-    found_types += {key for key in parser.defs["unions"]}
-    found_types += {key for key in parser.defs["enums"]}
-    found_functions += {key for key in parser.defs["functions"]}
-    found_values += {key for key in parser.defs["values"]}
-
-    for key, value in parser.defs["structs"].items():
-        struct_list[key] = Struct(key, value["members"])
-    for key, value in parser.defs["unions"].items():
-        struct_list[key] = Struct(key, value["members"])
-
-    for key, value in struct_list.items():
-        if key.startswith("anon_union") or key.startswith("anon_struct"):
-            continue
-
-        found_struct += [key]
-        discovered = value.discoverMembers(struct_list, key)
-        if discovered:
-            found_struct += discovered
-
-if len(found_functions) == 0:
-    raise RuntimeError(f'Parser found no functions. Is CUDA_HOME setup correctly? (CUDA_HOME="{CUDA_HOME}")')
+            found_struct += [key]
+            discovered = value.discoverMembers(struct_list, key)
+            if discovered:
+                found_struct += discovered
+
+    return found_types, found_functions, found_values, found_struct, struct_list
+
+
+include_path_list = [os.path.join(path, "include") for path in CUDA_HOME]
+header_dict = fetch_header_paths(required_headers, include_path_list)
+found_types, found_functions, found_values, found_struct, struct_list = parse_headers(header_dict)
 
 # ----------------------------------------------------------------------
 # Generate

cuda_core/cuda/core/experimental/__init__.py

@@ -4,7 +4,7 @@
 
 from cuda.core.experimental import utils
 from cuda.core.experimental._device import Device
-from cuda.core.experimental._event import EventOptions
+from cuda.core.experimental._event import Event, EventOptions
 from cuda.core.experimental._launcher import LaunchConfig, launch
 from cuda.core.experimental._linker import Linker, LinkerOptions
 from cuda.core.experimental._module import ObjectCode

cuda_core/cuda/core/experimental/_linker.py

@@ -349,11 +349,12 @@ class Linker:
     """
 
     class _MembersNeededForFinalize:
-        __slots__ = ("handle", "use_nvjitlink")
+        __slots__ = ("handle", "use_nvjitlink", "const_char_keep_alive")
 
         def __init__(self, program_obj, handle, use_nvjitlink):
             self.handle = handle
             self.use_nvjitlink = use_nvjitlink
+            self.const_char_keep_alive = []
             weakref.finalize(program_obj, self.close)
 
         def close(self):
@@ -390,27 +391,30 @@ def _add_code_object(self, object_code: ObjectCode):
         data = object_code._module
         assert_type(data, bytes)
         with _exception_manager(self):
+            name_str = f"{object_code._handle}_{object_code._code_type}"
             if _nvjitlink:
                 _nvjitlink.add_data(
                     self._mnff.handle,
                     self._input_type_from_code_type(object_code._code_type),
                     data,
                     len(data),
-                    f"{object_code._handle}_{object_code._code_type}",
+                    name_str,
                 )
             else:
+                name_bytes = name_str.encode()
                 handle_return(
                     _driver.cuLinkAddData(
                         self._mnff.handle,
                         self._input_type_from_code_type(object_code._code_type),
                         data,
                         len(data),
-                        f"{object_code._handle}_{object_code._code_type}".encode(),
+                        name_bytes,
                         0,
                         None,
                         None,
                     )
                 )
+                self._mnff.const_char_keep_alive.append(name_bytes)
 
     def link(self, target_type) -> ObjectCode:
         """
@@ -465,7 +469,7 @@ def get_error_log(self) -> str:
             _nvjitlink.get_error_log(self._mnff.handle, log)
         else:
             log = self._options.formatted_options[2]
-        return log.decode()
+        return log.decode("utf-8", errors="backslashreplace")
 
     def get_info_log(self) -> str:
         """Get the info log generated by the linker.
@@ -481,7 +485,7 @@ def get_info_log(self) -> str:
             _nvjitlink.get_info_log(self._mnff.handle, log)
         else:
             log = self._options.formatted_options[0]
-        return log.decode()
+        return log.decode("utf-8", errors="backslashreplace")
 
     def _input_type_from_code_type(self, code_type: str):
         # this list is based on the supported values for code_type in the ObjectCode class definition.

cuda_core/cuda/core/experimental/_program.py

@@ -503,7 +503,7 @@ def compile(self, target_type, name_expressions=(), logs=None):
                 if logsize > 1:
                     log = b" " * logsize
                     handle_return(nvrtc.nvrtcGetProgramLog(self._mnff.handle, log), handle=self._mnff.handle)
-                    logs.write(log.decode())
+                    logs.write(log.decode("utf-8", errors="backslashreplace"))
 
             return ObjectCode._init(data, target_type, symbol_mapping=symbol_mapping)
 

cuda_core/cuda/core/experimental/_utils/cuda_utils.py

@@ -95,7 +95,7 @@ def _check_error(error, handle=None):
             _, logsize = nvrtc.nvrtcGetProgramLogSize(handle)
             log = b" " * logsize
             _ = nvrtc.nvrtcGetProgramLog(handle, log)
-            err += f", compilation log:\n\n{log.decode()}"
+            err += f", compilation log:\n\n{log.decode('utf-8', errors='backslashreplace')}"
         raise NVRTCError(err)
     else:
         raise RuntimeError(f"Unknown error type: {error}")
@@ -203,6 +203,7 @@ def is_nested_sequence(obj):
     return is_sequence(obj) and any(is_sequence(elem) for elem in obj)
 
 
+@functools.lru_cache
 def get_binding_version():
     try:
         major_minor = importlib.metadata.version("cuda-bindings").split(".")[:2]

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

@@ -3,7 +3,7 @@
 ``cuda.core`` 0.2.0 Release Notes
 =================================
 
-Released on <TODO>, 2025
+Released on March 17, 2025
 
 Highlights
 ----------
@@ -27,10 +27,27 @@ New features
 - Expose :class:`ObjectCode` as a public API, which allows loading cubins from memory or disk. For loading other kinds of code types, please continue using :class:`Program`.
 - A C++ helper function ``get_cuda_native_handle()`` is provided in the new ``include/utility.cuh`` header to retrive the underlying CUDA C objects (ex: ``CUstream``) from a Python object returned by the ``.handle`` attribute (ex: :attr:`Stream.handle`).
 - For objects such as :class:`Program` and :class:`Linker` that could dispatch to different backends, a new ``.backend`` attribute is provided to query this information.
-- Support CUDA event timing.
-- An :class:`~_event.Event` may now be created without recording it to a :class:`~_stream.Stream` using the :meth:`Device.create_event` method.
+- Support CUDA :class:`Event` timing. (#481, #498, #508)
+- An :class:`Event` may now be created without recording it to a :class:`~_stream.Stream` using the :meth:`Device.create_event` method.
+- :class:`Program` now supports the additional ``PTX`` code type. (#317)
+- :meth:`Linker.link` exceptions now include the original error log. (#423)
+- In a systematic sweep through the cuda.core implementations, many exceptions messages were made more consistent and informative. (#458)
 
-Limitations
------------
-
-- <TODO>
+New examples
+------------
+- ``jit_lto_fractal.py`` — Demonstrates just-in-time link-time optimization for fractal generation. (:class:`Device`, :class:`LaunchConfig`, :class:`Linker`, :class:`LinkerOptions`, :class:`Program`, :class:`ProgramOptions`) (#475)
+- ``simple_multi_gpu_example.py`` — Example of using multiple GPUs. (:class:`Device`, :class:`Program`, :class:`LaunchConfig`) (#304)
+- ``show_device_properties.py`` — Displays detailed device properties. (:class:`Device`) (#474)
+
+Minor fixes and enhancements
+----------------------------
+- A dangling pointer problem in ``_linker.py`` was fixed. (#516)
+- Add ``@functools.lru_cache`` decorator for :func:`get_binding_version`. (#512)
+- Selected ``.decode()`` were changed to ``.decode("utf-8", errors="backslashreplace")`` to ensure that decoding error messages does not abort the process. (#510)
+- The performance of :meth:`Device.compute_capability` was improved. (#459)
+- The :class:`Program` constructor now issues a warning when falling back to :func:`cuLink`. (#315)
+- To avoid deprecation warnings, the cuda.bindings imports in the cuda.core implementations were cleaned up. (#404)
+
+Test fixes
+----------
+- Clean up device initialization in some tests. (#507)