Skip to content

fix(core): stop run_nvrtc corrupting the interned bytes singleton#15887

Open
harjothkhara wants to merge 1 commit into
NVIDIA-NeMo:mainfrom
harjothkhara:fix/run-nvrtc-bytes-singleton
Open

fix(core): stop run_nvrtc corrupting the interned bytes singleton#15887
harjothkhara wants to merge 1 commit into
NVIDIA-NeMo:mainfrom
harjothkhara:fix/run-nvrtc-bytes-singleton

Conversation

@harjothkhara

@harjothkhara harjothkhara commented Jul 7, 2026

Copy link
Copy Markdown

What does this PR do ?

Fixes a process-wide memory-corruption bug in run_nvrtc(): it allocated NVRTC output buffers as immutable b" " * size, and for an empty compile log (size == 1) CPython returns the interned, process-wide 1-byte singleton b" ". NVRTC then writes its C-string NUL terminator into that shared object, turning every b" " / bytes([32]) / 1-byte space slice in the whole process into b"\x00" permanently. Fix: use a mutable bytearray.

Collection: core (nemo/core/utils/cuda_python_utils.py)

Changelog

  • nemo/core/utils/cuda_python_utils.py: allocate the NVRTC log buffer and PTX buffer as bytearray(size) / bytearray(ptxSize) instead of b" " * size (mutable, never interned).
  • Convert with bytes(ptx) before np.char.array(...) so downstream behavior is byte-for-byte identical (passing a bytearray directly to np.char.array would be misinterpreted as a sequence of ints).
  • tests/core/test_cuda_python_utils.py (new): GPU-free regression test that fakes the NVRTC/CUDA bindings and emulates NVRTC writing into the buffer via memoryview; asserts the buffers are mutable and the interned singleton is intact. Fails on the old code (TypeError: cannot modify read-only memory), passes on the fix. Skips when cuda-python is unavailable, matching tests/collections/asr/decoding/test_cuda_graph_rnnt_greedy_decoding.py.

Usage

Reproduction (self-contained, no GPU / no NeMo import — just cuda-python), from #15790:

from cuda.bindings import nvrtc
err, prog = nvrtc.nvrtcCreateProgram(b'extern "C" __global__ void k(){}\n', b"k.cu", 0, [], [])
(err,) = nvrtc.nvrtcCompileProgram(prog, 0, [])
err, size = nvrtc.nvrtcGetProgramLogSize(prog)   # == 1 for a clean compile
buf = b" " * size                                # the interned singleton b" "
(err,) = nvrtc.nvrtcGetProgramLog(prog, buf)     # NVRTC writes NUL into it
print(bytes([32]))                               # -> b'\x00'  (b" " is now corrupted process-wide)

GitHub Actions CI

CI has not been triggered yet (external fork). A maintainer can start it with /ok to test <sha> / "Approve and run" and the Run CICD label.

Before your PR is "Ready for review"

Pre checks:

  • Make sure you read and followed Contributor guidelines
  • Did you write any new necessary tests? — added tests/core/test_cuda_python_utils.py
  • Did you add or update any necessary documentation? — n/a (internal bugfix)
  • Does the PR affect components that are optional to install? (cuda-python)
    • Reviewer: import guards are unchanged — run_nvrtc stays behind @cuda_python_required, and the new test pytest.skips when cuda-python is unavailable.

PR Type:

  • New Feature
  • Bugfix
  • Documentation

Additional Information

run_nvrtc allocated its NVRTC output buffers as `b" " * size`. NVRTC
writes the compile log / PTX (a C string, including its NUL terminator)
into those buffers, but `bytes` is immutable. For an empty compile log
`size == 1`, and CPython's `b" " * 1` returns the interned, process-wide
1-byte singleton `b" "`; NVRTC's NUL write then turns every `b" "`,
`bytes([32])` and 1-byte space slice in the whole process into `b"\x00"`
permanently.

Allocate the log and PTX buffers as `bytearray(size)` (mutable and never
interned) and convert with `bytes(ptx)` before `np.char.array` to keep
the existing behavior. Adds a GPU-free regression test that fakes the
NVRTC/CUDA bindings and asserts the output buffers are mutable and the
interned singleton is left intact.

Fixes NVIDIA-NeMo#15790

Signed-off-by: harjoth <harjoth.khara@gmail.com>
@copy-pr-bot

copy-pr-bot Bot commented Jul 7, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@github-actions github-actions Bot added core Changes to NeMo Core community-request labels Jul 7, 2026
@nithinraok nithinraok requested a review from artbataev July 7, 2026 19:14
@svcnvidia-nemo-ci svcnvidia-nemo-ci added the waiting-on-maintainers Waiting on maintainers to respond label Jul 9, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

community-request core Changes to NeMo Core waiting-on-maintainers Waiting on maintainers to respond

Projects

None yet

Development

Successfully merging this pull request may close these issues.

run_nvrtc() corrupts CPython's interned 1-byte bytes singleton b' ' when the NVRTC compile log is empty

2 participants