From 1ae478bc8c61955930d26699682cf26db6dbd121 Mon Sep 17 00:00:00 2001 From: Rui Luo Date: Fri, 16 Oct 2026 16:37:02 +0800 Subject: [PATCH] fix(cuda.core): normalize PCH status after heap retry Convert the retry status to str before storing it in the Cython str field, matching the initial compilation path and preserving the public PCHStatusType result. Add a regression that forces PCH heap growth and restores the original heap size, and skip it when NVRTC_DISABLE_PCH is set. Signed-off-by: Rui Luo --- cuda_core/cuda/core/_program.pyx | 2 +- cuda_core/docs/source/release/1.3.0-notes.rst | 3 ++ cuda_core/tests/test_program.py | 31 +++++++++++++++++++ 3 files changed, 35 insertions(+), 1 deletion(-) diff --git a/cuda_core/cuda/core/_program.pyx b/cuda_core/cuda/core/_program.pyx index c39b03d0cd7..94f5bf74f72 100644 --- a/cuda_core/cuda/core/_program.pyx +++ b/cuda_core/cuda/core/_program.pyx @@ -1100,7 +1100,7 @@ cdef object Program_compile_nvrtc(Program self, str target_type, object name_exp ) status = _read_pch_status(retry_prog) - self._pch_status = status if status is not None else str(PCHStatusType.FAILED) + self._pch_status = str(status) if status is not None else str(PCHStatusType.FAILED) return result diff --git a/cuda_core/docs/source/release/1.3.0-notes.rst b/cuda_core/docs/source/release/1.3.0-notes.rst index f292bb209a6..648906f84bb 100644 --- a/cuda_core/docs/source/release/1.3.0-notes.rst +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -104,6 +104,9 @@ New features Fixes and enhancements ---------------------- +- Fixed a ``TypeError`` in :meth:`Program.compile` when retrying after PCH + heap exhaustion. + - :attr:`system.Device.uuid_without_prefix` preserves UUIDs that NVML already reports without a prefix, such as on Orin. This also allows :meth:`system.Device.to_cuda_device` to match those devices correctly. diff --git a/cuda_core/tests/test_program.py b/cuda_core/tests/test_program.py index 256b3553173..b51554a1603 100644 --- a/cuda_core/tests/test_program.py +++ b/cuda_core/tests/test_program.py @@ -298,6 +298,37 @@ def test_cpp_program_pch_auto_creates(init_cuda, tmp_path): program.close() +@nvrtc_pch_available +@pytest.mark.skipif("NVRTC_DISABLE_PCH" in os.environ, reason="PCH is disabled by NVRTC_DISABLE_PCH") +@pytest.mark.thread_unsafe(reason="resizes the process-global NVRTC PCH heap") +@pytest.mark.agent_authored(model="gpt-6") +def test_cpp_program_pch_heap_retry(tmp_path): + """PCH heap exhaustion retries compilation and returns a typed creation status.""" + source = '\n#pragma nv_hdrstop\nextern "C" __global__ void my_kernel() {}\n' + pch_path = tmp_path / "retry.pch" + arch = min(handle_return(nvrtc.nvrtcGetSupportedArchs())) + options = ProgramOptions(arch=f"compute_{arch}", name="retry.cu", create_pch=str(pch_path)) + previous_heap = handle_return(nvrtc.nvrtcGetPCHHeapSize()) + try: + # NVRTC rounds this up; compare against the actual initial heap size. + handle_return(nvrtc.nvrtcSetPCHHeapSize(1)) + small_heap = handle_return(nvrtc.nvrtcGetPCHHeapSize()) + program = Program(source, "c++", options) + try: + result = program.compile("ptx") + assert isinstance(result, ObjectCode) + assert b"my_kernel" in result.code + assert program.pch_status is PCHStatusType.CREATED + assert pch_path.is_file() + # Success alone could miss the retry path; require automatic growth. + assert handle_return(nvrtc.nvrtcGetPCHHeapSize()) > small_heap + finally: + program.close() + finally: + handle_return(nvrtc.nvrtcSetPCHHeapSize(previous_heap)) + assert handle_return(nvrtc.nvrtcGetPCHHeapSize()) == previous_heap + + @bundled_headers_available @pytest.mark.agent_authored(model="claude-sonnet-5") def test_use_bundled_headers_installs_and_compiles(init_cuda, tmp_path, monkeypatch):