Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion cuda_core/cuda/core/_program.pyx
Original file line number Diff line number Diff line change
Expand Up @@ -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


Expand Down
3 changes: 3 additions & 0 deletions cuda_core/docs/source/release/1.3.0-notes.rst
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
31 changes: 31 additions & 0 deletions cuda_core/tests/test_program.py
Original file line number Diff line number Diff line change
Expand Up @@ -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):
Expand Down
Loading