Fix #2363: Support using bundled headers with NVRTC 13.3+.

cuda_core/cuda/core/_program.pyi

@@ -218,6 +218,13 @@ class ProgramOptions:
     include_path : str | list[str], optional
         Add the directory or directories to the list of directories to be searched for headers.
         Default: None
+    use_bundled_headers : bool, optional
+        Use the CUDA and CCCL headers bundled with NVRTC, installed into a per-user cache
+        directory, instead of requiring a full CUDA Toolkit installation. Implemented via NVRTC's
+        ``--use-bundled-headers=<dir>`` compiler option, which installs the headers into the cache
+        directory (skipping installation if already present and up to date) and adds that
+        directory to the include search path. NVRTC only.
+        Default: False
     pre_include : str | list[str], optional
         Preinclude one or more headers during preprocessing. Can be either a string or a list of strings.
         Default: None
@@ -350,6 +357,7 @@ class ProgramOptions:
     define_macro: str | tuple[str, str] | list[str | tuple[str, str]] | tuple[str | tuple[str, str], ...] | None = None
     undefine_macro: str | list[str] | tuple[str] | None = None
     include_path: str | list[str] | tuple[str] | None = None
+    use_bundled_headers: bool | None = None
     pre_include: str | list[str] | tuple[str] | None = None
     no_source_include: bool | None = None
     std: str | None = None

cuda_core/cuda/core/_program.pyx

@@ -47,6 +47,7 @@ from cuda.core._utils.cuda_utils import (
     is_sequence,
 )
 from cuda.core._utils.version import binding_version, driver_version
+from cuda.core.utils._cache_dir import _default_cache_dir
 from cuda.core.typing import ObjectCodeFormatType, CompilerBackendType, PCHStatusType, SourceCodeType
 
 __all__ = ["Program", "ProgramOptions"]
@@ -430,6 +431,13 @@ class ProgramOptions:
     include_path : str | list[str], optional
         Add the directory or directories to the list of directories to be searched for headers.
         Default: None
+    use_bundled_headers : bool, optional
+        Use the CUDA and CCCL headers bundled with NVRTC, installed into a per-user cache
+        directory, instead of requiring a full CUDA Toolkit installation. Implemented via NVRTC's
+        ``--use-bundled-headers=<dir>`` compiler option, which installs the headers into the cache
+        directory (skipping installation if already present and up to date) and adds that
+        directory to the include search path. NVRTC only.
+        Default: False
     pre_include : str | list[str], optional
         Preinclude one or more headers during preprocessing. Can be either a string or a list of strings.
         Default: None
@@ -563,6 +571,7 @@ class ProgramOptions:
     define_macro: str | tuple[str, str] | list[str | tuple[str, str]] | tuple[str | tuple[str, str], ...] | None = None
     undefine_macro: str | list[str] | tuple[str] | None = None
     include_path: str | list[str] | tuple[str] | None = None
+    use_bundled_headers: bool | None = None
     pre_include: str | list[str] | tuple[str] | None = None
     no_source_include: bool | None = None
     std: str | None = None
@@ -606,6 +615,16 @@ class ProgramOptions:
         # Set arch to default if not provided
         if self.arch is None:
             self.arch = f"sm_{Device().arch}"
+        if self.use_bundled_headers:
+            # --use-bundled-headers (and the bundled CUDA/CCCL headers themselves) were
+            # introduced in NVRTC 13.3.
+            nvrtc_major, nvrtc_minor = handle_return(nvrtc.nvrtcVersion())
+            if (nvrtc_major, nvrtc_minor) < (13, 3):
+                raise RuntimeError(
+                    "use_bundled_headers requires NVRTC >= 13.3, but found "
+                    f"{nvrtc_major}.{nvrtc_minor}. Upgrade the CUDA Toolkit / driver providing "
+                    "libnvrtc, or set use_bundled_headers=False and supply include_path manually."
+                )
         if self.extra_sources is not None:
             if not is_sequence(self.extra_sources):
                 raise TypeError(
@@ -765,8 +784,6 @@ def _find_libdevice_path() -> object:
     return find_bitcode_lib("device")
 
 
-
-
 cdef inline bint _process_define_macro_inner(list options, object macro) except? -1:
     """Process a single define macro, returning True if successful."""
     if isinstance(macro, str):
@@ -1234,6 +1251,8 @@ cdef inline list _prepare_nvrtc_options_impl(object opts):
         elif is_sequence(opts.undefine_macro):
             for macro in opts.undefine_macro:
                 options.append(f"--undefine-macro={macro}")
+    if opts.use_bundled_headers:
+        options.append(f"--use-bundled-headers={_default_cache_dir() / 'nvrtc-bundled-headers'}")
     if opts.include_path is not None:
         if isinstance(opts.include_path, str):
             options.append(f"--include-path={opts.include_path}")
@@ -1398,6 +1417,8 @@ cdef inline object _prepare_nvvm_options_impl(object opts, bint as_bytes):
         unsupported.append("undefine_macro")
     if opts.include_path is not None:
         unsupported.append("include_path")
+    if opts.use_bundled_headers:
+        unsupported.append("use_bundled_headers")
     if opts.pre_include is not None:
         unsupported.append("pre_include")
     if opts.no_source_include is not None and opts.no_source_include:

cuda_core/cuda/core/utils/_cache_dir.py

@@ -0,0 +1,41 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Shared user-cache-root resolution for cuda.core's on-disk caches."""
+
+from __future__ import annotations
+
+import os
+from pathlib import Path
+
+# Exposed as a module-level flag so tests can toggle it without monkeypatching
+# ``os.name`` itself (pathlib reads ``os.name`` at instantiation time).
+_IS_WINDOWS = os.name == "nt"
+
+
+def _default_cache_dir() -> Path:
+    """OS-conventional root for cuda.core's on-disk caches.
+
+    Resolves to the user-cache root for the calling user, with a
+    ``cuda-python`` vendor leaf so callers can each place their own cache
+    under a stable, shared root:
+
+    * Linux: ``$XDG_CACHE_HOME/cuda-python``
+      (default ``~/.cache/cuda-python`` per the XDG Base Directory spec).
+    * Windows: ``%LOCALAPPDATA%\\cuda-python``
+      (Windows uses local AppData -- caches don't roam; falls back to
+      ``~/AppData/Local`` if the env var is unset).
+
+    CUDA does not support macOS, so no macOS branch is provided.
+
+    Callers append their own leaf directory, e.g. ``program-cache`` or
+    ``nvrtc-headers``.
+    """
+    if _IS_WINDOWS:
+        local_app_data = os.environ.get("LOCALAPPDATA")
+        root = Path(local_app_data) if local_app_data else Path.home() / "AppData" / "Local"
+    else:
+        xdg = os.environ.get("XDG_CACHE_HOME")
+        root = Path(xdg) if xdg else Path.home() / ".cache"
+    return root / "cuda-python"

cuda_core/cuda/core/utils/_program_cache/_file_stream.py

@@ -23,6 +23,7 @@
 from typing import Any, Callable, Iterable
 
 from cuda.core._module import ObjectCode
+from cuda.core.utils._cache_dir import _default_cache_dir as _user_cache_dir
 
 from ._abc import ProgramCacheResource, _as_key_bytes, _extract_bytes
 
@@ -57,28 +58,10 @@ def _stat_key(st: os.stat_result) -> tuple[int, int, int]:
 
 
 def _default_cache_dir() -> Path:
-    """OS-conventional default location for the file-stream cache.
-
-    Resolves to the user-cache root for the calling user, with a
-    ``program-cache`` leaf so future tooling can place sibling caches
-    under the same ``cuda-python`` vendor directory:
-
-    * Linux: ``$XDG_CACHE_HOME/cuda-python/program-cache``
-      (default ``~/.cache/cuda-python/program-cache`` per the XDG Base
-      Directory spec).
-    * Windows: ``%LOCALAPPDATA%\\cuda-python\\program-cache``
-      (Windows uses local AppData -- caches don't roam; falls back to
-      ``~/AppData/Local`` if the env var is unset).
-
-    CUDA does not support macOS, so no macOS branch is provided.
+    """Default location for the file-stream cache: the ``program-cache`` leaf under the shared
+    ``cuda-python`` user-cache root (see :func:`cuda.core.utils._cache_dir._default_cache_dir`).
     """
-    if _IS_WINDOWS:
-        local_app_data = os.environ.get("LOCALAPPDATA")
-        root = Path(local_app_data) if local_app_data else Path.home() / "AppData" / "Local"
-    else:
-        xdg = os.environ.get("XDG_CACHE_HOME")
-        root = Path(xdg) if xdg else Path.home() / ".cache"
-    return root / "cuda-python" / "program-cache"
+    return _user_cache_dir() / "program-cache"
 
 
 def _with_sharing_retry(

cuda_core/tests/test_cache_dir.py

@@ -0,0 +1,44 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import pytest
+
+
+@pytest.mark.agent_authored(model="claude-sonnet-5")
+def test_user_cache_dir_lives_under_platform_root(monkeypatch, tmp_path):
+    """The shared user-cache root (``cuda.core.utils._cache_dir``) is platform-specific:
+
+    * Linux: ``$XDG_CACHE_HOME`` or ``~/.cache``.
+    * Windows: ``%LOCALAPPDATA%`` or ``~/AppData/Local``.
+
+    Both branches must end in ``cuda-python``; that suffix is what guarantees a
+    stable on-disk layout across releases, and callers (e.g. the file-stream
+    cache, NVRTC's bundled-headers cache) each append their own leaf under it.
+    """
+    from pathlib import Path
+
+    from cuda.core.utils import _cache_dir
+    from cuda.core.utils._cache_dir import _default_cache_dir
+
+    # Path must end with cuda-python regardless of platform.
+    assert _default_cache_dir().parts[-1] == "cuda-python"
+
+    # Linux branch: XDG_CACHE_HOME wins when set.
+    monkeypatch.setattr(_cache_dir, "_IS_WINDOWS", False)
+    monkeypatch.setenv("XDG_CACHE_HOME", str(tmp_path / "xdg"))
+    assert _default_cache_dir() == tmp_path / "xdg" / "cuda-python"
+
+    # Linux branch: falls back to ``~/.cache`` when XDG_CACHE_HOME is unset.
+    monkeypatch.delenv("XDG_CACHE_HOME", raising=False)
+    monkeypatch.setattr(Path, "home", classmethod(lambda _cls: tmp_path / "home"))
+    assert _default_cache_dir() == tmp_path / "home" / ".cache" / "cuda-python"
+
+    # Windows branch: LOCALAPPDATA wins when set.
+    monkeypatch.setattr(_cache_dir, "_IS_WINDOWS", True)
+    monkeypatch.setenv("LOCALAPPDATA", str(tmp_path / "appdata"))
+    assert _default_cache_dir() == tmp_path / "appdata" / "cuda-python"
+
+    # Windows branch: falls back to ``~/AppData/Local`` when LOCALAPPDATA is unset.
+    monkeypatch.delenv("LOCALAPPDATA", raising=False)
+    assert _default_cache_dir() == tmp_path / "home" / "AppData" / "Local" / "cuda-python"

cuda_core/tests/test_program.py

@@ -72,6 +72,11 @@ def _has_nvrtc_pch_apis_for_tests():
     reason="PCH runtime APIs require NVRTC >= 12.8 bindings",
 )
 
+bundled_headers_available = pytest.mark.skipif(
+    (_get_nvrtc_version_for_tests() or 0) < 13300,
+    reason="use_bundled_headers requires NVRTC >= 13.3",
+)
+
 
 def _has_check_nvvm_compiler_options():
     try:
@@ -302,6 +307,48 @@ def test_cpp_program_pch_auto_creates(init_cuda, tmp_path):
     program.close()
 
 
+@bundled_headers_available
+@pytest.mark.agent_authored(model="claude-sonnet-5")
+def test_use_bundled_headers_installs_and_compiles(init_cuda, tmp_path, monkeypatch):
+    """``use_bundled_headers`` should install NVRTC's bundled CUDA/CCCL headers into the
+    (monkeypatched) cache directory and make them available on the include path, without
+    a CUDA Toolkit or any user-supplied ``include_path``."""
+    import cuda.core._program as _program_module
+
+    cache_root = tmp_path / "cache-root"
+    monkeypatch.setattr(_program_module, "_default_cache_dir", lambda: cache_root)
+
+    code = """
+#include <cuda/std/type_traits>
+extern "C" __global__ void my_kernel(int *out) {
+    *out = cuda::std::is_integral<int>::value;
+}
+"""
+    headers_dir = cache_root / "nvrtc-bundled-headers"
+    assert not headers_dir.exists()
+
+    # Sanity check: without use_bundled_headers, the CCCL header isn't found (proves the
+    # option -- not some ambient CUDA Toolkit install -- is what makes the compile below work).
+    program = Program(code, "c++")
+    try:
+        with pytest.raises(CUDAError, match="could not open source file"):
+            program.compile("ptx")
+    finally:
+        program.close()
+
+    program = Program(code, "c++", ProgramOptions(use_bundled_headers=True))
+    try:
+        object_code = program.compile("ptx")
+    finally:
+        program.close()
+    assert isinstance(object_code, ObjectCode)
+
+    assert headers_dir.is_dir()
+    assert (headers_dir / ".nvrtc_headers_version").is_file()
+    assert (headers_dir / "cccl").is_dir()
+    assert (headers_dir / "cccl" / "cuda" / "std" / "type_traits").is_file()
+
+
 def test_cpp_program_pch_status_none_without_pch(init_cuda):
     code = 'extern "C" __global__ void my_kernel() {}'
     program = Program(code, "c++")

cuda_core/tests/test_program_cache.py

@@ -1526,41 +1526,18 @@ def test_filestream_cache_rejects_non_positive_size_cap(tmp_path, bad):
         FileStreamProgramCache(tmp_path / "fc", max_size_bytes=bad)
 
 
-def test_default_cache_dir_lives_under_user_cache_root(monkeypatch, tmp_path):
-    """The cache root is platform-specific:
-
-    * Linux: ``$XDG_CACHE_HOME`` or ``~/.cache``.
-    * Windows: ``%LOCALAPPDATA%`` or ``~/AppData/Local``.
-
-    Both branches must end in ``cuda-python/program-cache``; that suffix
-    is what guarantees a stable on-disk layout across releases.
-    """
-    from pathlib import Path
-
+@pytest.mark.agent_authored(model="claude-sonnet-5")
+def test_file_stream_default_cache_dir_appends_program_cache_leaf(monkeypatch, tmp_path):
+    """The file-stream cache's default dir is the shared user-cache root plus a
+    ``program-cache`` leaf, so it can live alongside sibling caches (e.g. NVRTC's
+    bundled-headers cache) under the same ``cuda-python`` vendor directory."""
     from cuda.core.utils import _program_cache
-    from cuda.core.utils._program_cache._file_stream import _default_cache_dir
 
-    # Path must end with cuda-python/program-cache regardless of platform.
-    assert _default_cache_dir().parts[-2:] == ("cuda-python", "program-cache")
+    monkeypatch.setattr(_program_cache._file_stream, "_user_cache_dir", lambda: tmp_path / "root")
 
-    # Linux branch: XDG_CACHE_HOME wins when set.
-    monkeypatch.setattr(_program_cache._file_stream, "_IS_WINDOWS", False)
-    monkeypatch.setenv("XDG_CACHE_HOME", str(tmp_path / "xdg"))
-    assert _default_cache_dir() == tmp_path / "xdg" / "cuda-python" / "program-cache"
-
-    # Linux branch: falls back to ``~/.cache`` when XDG_CACHE_HOME is unset.
-    monkeypatch.delenv("XDG_CACHE_HOME", raising=False)
-    monkeypatch.setattr(Path, "home", classmethod(lambda _cls: tmp_path / "home"))
-    assert _default_cache_dir() == tmp_path / "home" / ".cache" / "cuda-python" / "program-cache"
-
-    # Windows branch: LOCALAPPDATA wins when set.
-    monkeypatch.setattr(_program_cache._file_stream, "_IS_WINDOWS", True)
-    monkeypatch.setenv("LOCALAPPDATA", str(tmp_path / "appdata"))
-    assert _default_cache_dir() == tmp_path / "appdata" / "cuda-python" / "program-cache"
+    from cuda.core.utils._program_cache._file_stream import _default_cache_dir
 
-    # Windows branch: falls back to ``~/AppData/Local`` when LOCALAPPDATA is unset.
-    monkeypatch.delenv("LOCALAPPDATA", raising=False)
-    assert _default_cache_dir() == tmp_path / "home" / "AppData" / "Local" / "cuda-python" / "program-cache"
+    assert _default_cache_dir() == tmp_path / "root" / "program-cache"
 
 
 def test_filestream_cache_uses_default_dir_when_path_omitted(tmp_path, monkeypatch):