diff --git a/cuda_core/cuda/core/_program.pyi b/cuda_core/cuda/core/_program.pyi index f739b7662a9..ca18850e3b5 100644 --- a/cuda_core/cuda/core/_program.pyi +++ b/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=`` 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 diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index b937ded7548..9ea624f8e83 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/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=`` 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: diff --git a/cuda_core/cuda/core/utils/_cache_dir.py b/cuda_core/cuda/core/utils/_cache_dir.py new file mode 100644 index 00000000000..2981db14da3 --- /dev/null +++ b/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" diff --git a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py b/cuda_core/cuda/core/utils/_program_cache/_file_stream.py index eb71abf5446..314c7b612bc 100644 --- a/cuda_core/cuda/core/utils/_program_cache/_file_stream.py +++ b/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( diff --git a/cuda_core/tests/test_cache_dir.py b/cuda_core/tests/test_cache_dir.py new file mode 100644 index 00000000000..7d431181ff4 --- /dev/null +++ b/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" diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 4874c6184c7..e457b437fc8 100644 --- a/cuda_core/tests/test_program.py +++ b/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 +extern "C" __global__ void my_kernel(int *out) { + *out = cuda::std::is_integral::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++") diff --git a/cuda_core/tests/test_program_cache.py b/cuda_core/tests/test_program_cache.py index 2b5402374a6..bd327eb0d05 100644 --- a/cuda_core/tests/test_program_cache.py +++ b/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):