Repository navigation
Out of bound access in the test_graph_update test #843
Description
CI log: https://github.com/NVIDIA/cuda-python/actions/runs/16994945071/job/48183959433
tests/test_graph.py::test_graph_update ========= Invalid __global__ read of size 4 bytes
========= at add_one(int *)+0x40
========= by thread (0,0,0) in block (0,0,0)
========= Access at 0x7c547b000000 is out of bounds
========= and is inside the nearest allocation at 0x7c547b000000 of size 12 bytes
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x3f2bd] in cydriver.cpython-312-x86_64-linux-gnu.so
========= Host Frame: [0xf4e9] in cydriver.cpython-312-x86_64-linux-gnu.so
========= Host Frame: [0x3fe25d] in driver.cpython-312-x86_64-linux-gnu.so
========= Host Frame: PyObject_Vectorcall in call.c:325 [0x210c94] in libpython3.12.so.1.0
========= Host Frame: _PyEval_EvalFrameDefault in bytecodes.c:2715 [0x2606fd] in libpython3.12.so.1.0
========= Host Frame: _PyObject_Call_Prepend in call.c:508 [0x2114b6] in libpython3.12.so.1.0
========= Host Frame: slot_tp_call in typeobject.c:8791 [0x2d3572] in libpython3.12.so.1.0
========= Host Frame: _PyObject_MakeTpCall in call.c:240 [0x210962] in libpython3.12.so.1.0
========= Host Frame: _PyEval_EvalFrameDefault in bytecodes.c:2715 [0x2606fd] in libpython3.12.so.1.0
========= Host Frame: _PyObject_Call_Prepend in call.c:508 [0x2114b6] in libpython3.12.so.1.0
========= Host Frame: slot_tp_call in typeobject.c:8791 [0x2d3572] in libpython3.12.so.1.0
========= Host Frame: PyObject_Call [0x211b9e] in libpython3.12.so.1.0
========= Host Frame: _PyEval_EvalFrameDefault in bytecodes.c:3263 [0x264f1b] in libpython3.12.so.1.0
========= Host Frame: _PyObject_Call_Prepend in call.c:508 [0x2114b6] in libpython3.12.so.1.0
========= Host Frame: slot_tp_call in typeobject.c:8791 [0x2d3572] in libpython3.12.so.1.0
========= Host Frame: _PyObject_MakeTpCall in call.c:240 [0x210962] in libpython3.12.so.1.0
========= Host Frame: _PyEval_EvalFrameDefault in bytecodes.c:2715 [0x2606fd] in libpython3.12.so.1.0
========= Host Frame: _PyObject_Call_Prepend in call.c:508 [0x2114b6] in libpython3.12.so.1.0
========= Host Frame: slot_tp_call in typeobject.c:8791 [0x2d3572] in libpython3.12.so.1.0
========= Host Frame: _PyObject_MakeTpCall in call.c:240 [0x210962] in libpython3.12.so.1.0
========= Host Frame: _PyEval_EvalFrameDefault in bytecodes.c:2715 [0x2606fd] in libpython3.12.so.1.0
========= Host Frame: _PyObject_Call_Prepend in call.c:508 [0x2114b6] in libpython3.12.so.1.0
========= Host Frame: slot_tp_call in typeobject.c:8791 [0x2d3572] in libpython3.12.so.1.0
========= Host Frame: _PyObject_MakeTpCall in call.c:240 [0x210962] in libpython3.12.so.1.0
========= Host Frame: _PyEval_EvalFrameDefault in bytecodes.c:2715 [0x2606fd] in libpython3.12.so.1.0
========= Host Frame: PyEval_EvalCode in ceval.c:580 [0x2fa2f5] in libpython3.12.so.1.0
========= Host Frame: run_eval_code_obj in pythonrun.c:1757 [0x31b427] in libpython3.12.so.1.0
========= Host Frame: run_mod in pythonrun.c:1778 [0x31af6a] in libpython3.12.so.1.0
========= Host Frame: pyrun_file in pythonrun.c:1674 [0x31aeb0] in libpython3.12.so.1.0
========= Host Frame: _PyRun_SimpleFileObject in pythonrun.c:459 [0x31ace8] in libpython3.12.so.1.0
========= Host Frame: _PyRun_AnyFileObject in pythonrun.c:78 [0x31ab06] in libpython3.12.so.1.0
========= Host Frame: Py_RunMain in main.c:714 [0x324274] in libpython3.12.so.1.0
========= Host Frame: Py_BytesMain in main.c:768 [0x323d3c] in libpython3.12.so.1.0
========= Host Frame: [0x29d8f] in libc.so.6
========= Host Frame: __libc_start_main [0x29e3f] in libc.so.6
========= Host Frame: _start [0x1094] in python
========= Host Frame: launch in _graph.py:782
========= Host Frame: test_graph_update in test_graph.py:612
========= Host Frame: pytest_pyfunc_call in python.py:157
========= Host Frame: _multicall in _callers.py:121
========= Host Frame: _hookexec in _manager.py:120
========= Host Frame: __call__ in _hooks.py:512
========= Host Frame: runtest in python.py:1671
========= Host Frame: pytest_runtest_call in runner.py:178
========= Host Frame: _multicall in _callers.py:121
========= Host Frame: _hookexec in _manager.py:120
========= Host Frame: __call__ in _hooks.py:512
========= Host Frame: <lambda> in runner.py:246
========= Host Frame: from_call in runner.py:344
========= Host Frame: call_and_report in runner.py:245
========= Host Frame: runtestprotocol in runner.py:136
========= Host Frame: pytest_runtest_protocol in runner.py:117
========= Host Frame: _multicall in _callers.py:121
========= Host Frame: _hookexec in _manager.py:120
========= Host Frame: __call__ in _hooks.py:512
========= Host Frame: pytest_runtestloop in main.py:367
========= Host Frame: _multicall in _callers.py:121
========= Host Frame: _hookexec in _manager.py:120
========= Host Frame: __call__ in _hooks.py:512
========= Host Frame: _main in main.py:343
========= Host Frame: wrap_session in main.py:289
========= Host Frame: pytest_cmdline_main in main.py:336
========= Host Frame: _multicall in _callers.py:121
========= Host Frame: _hookexec in _manager.py:120
========= Host Frame: __call__ in _hooks.py:512
========= Host Frame: main in __init__.py:175
=========
FAILED [ 38%]
Metadata
Metadata
Assignees
Type
Projects
- Status
Activity
leofang commentedon Aug 15, 2025
This might explain the occasional CI crashes we've seen. @oleksandr-pavlyk would you have time next week to take a look?
oleksandr-pavlyk commentedon Aug 18, 2025
Can reproduce, sporadically.
Creating environment
Create and activate conda environment using the explicit file provided above.
Then, using compute sanitizer from 12.9 CTK,
/usr/local/cuda/bin/compute-sanitizer \ --target-processes=all \ --launch-timeout=0 \ --tool=memcheck \ --error-exitcode=1 \ --report-api-errors=no \ --padding=32 \ pytest repos/cuda-python/cuda_core/tests/It helps to remove/reinstate the padding CLI argument to expedite reproducing. The issue observed:
The traceback points to exception occurring while launching the first graph:
https://github.com/NVIDIA/cuda-python/blob/main/cuda_core/tests/test_graph.py#L607-L616
oleksandr-pavlyk commentedon Aug 18, 2025
Incidentally, if
test_graph.pyis being tested stand-alone, the failure occurs in a different test:failing in
test_graph.py::test_graph_conditional_switch[0]withcuda.core.experimental._utils.cuda_utils.CUDAError: CUDA_ERROR_ILLEGAL_ADDRESS.When
--report-api-errorsis set toall, even for green run of the test file, a lot of errors likeare reported by compute sanitizer.
oleksandr-pavlyk commentedon Aug 18, 2025
The
CUDA_ERROR_INVALID_CONTEXTarises frominit_cudafixture:Note the following host frames:
These may be unrelated to the out-of-bound access issue though
leofang commentedon Aug 18, 2025
btw cuda.core currently requires
--report-api-errors=nooleksandr-pavlyk commentedon Aug 18, 2025
No real progress, but the issue can be reproduced with
Reproducing stand-alone script
Stand-alone it runs fine:
Launching the same with compute sanitizer
/usr/local/cuda/bin/compute-sanitizer --report-api-errors=no python run.pytripsCUDAError: CUDA_ERROR_ILLEGAL_ADDRESS.oleksandr-pavlyk commentedon Aug 18, 2025
I should add that my driver is 575.57.08:
oleksandr-pavlyk commentedon Aug 18, 2025
@leofang I find that changing the example away from using
LegacyPinnedMemoryResourceto usingDeviceMemoryResourceand injecting copies from device allocation back to host for validation resolves the crash.I can open a PR to see if this would fix the issue with CI. Pursuing triage with pinned memory separately.
kkraus14 commentedon Aug 19, 2025
It sounds like we're missing a synchronization somewhere then where running against pinned host memory would keep things asynchronous, whereas presumably the device to host copy you injected used pageable host memory which synchronized things?
leofang commentedon Aug 19, 2025
Sasha mentioned offline that the sanitizer team noted this might be a compute-sanitizer bug. He'll update this thread soon-ish after investigation.
rwgk commentedon Aug 20, 2025
JIC it's useful, here is another flake of this kind:
The rerun succeeded.
oleksandr-pavlyk commentedon Aug 21, 2025
I have finished building C++ stand-alone executable that builds same graphs as
test_graph.py. It runs fine both stand-alone, and under CTK 12.9 compute sanitizer:I have attempted several dozen runs, without running into any problem.
To execute ELF, two files are expected to be present in the same folder,
"simple_kernels.cubin"and"conditional_kernels.cubin". These were constructed as follows:test_graph.cpp
Note that this source file won't compile for CTK 13.0, due to API changes in
cuStreamGetCaptureInfobetween 12.9 and 13.0 as well as related changes to other graph APIs.Fix an apparent mistake in GraphBuilder.add_child
Fix an apparent mistake in GraphBuilder.add_child (#879)