CPU-only pipeline crashes on second feed_input due to cudaGetDevice() in AccessOrder::wait() host branch

[fversaci]

Version

2.1

Describe the bug.

I've been experiencing some crashes, with cuda errors, when feeding, more than once, data into a cpu-only pipeline:

pipe.feed_input("Reader", batch1)
results1 = pipe.run()      # works
pipe.feed_input("Reader", batch2)  # ← crash here
results2 = pipe.run()

I had minimax 2.7 trace this crash and it seems to be originated by a bug in AccessOrder::wait():

A device-ordered AccessOrder leaks into a CPU-only pipeline when a GPU conversion (.gpu()) is chained downstream. On the second feed_input, AccessOrder::wait() is called with the this side in host order and other in device order, which enters the host branch and unconditionally calls cudaGetDevice(). On systems with CUDA installed, cudaStreamSynchronize() is then called on a stale device stream handle. The result is non-deterministic: a RuntimeError, silent segfault, or hang.

Minimum reproducible example

Here's a repo which minimally reproduces the crash:
https://github.com/fversaci/feed-bug

Relevant log output

$ uv run main.py
Loaded plugin: /mnt/tdm-dic/users/cesco-safe/code/feed-bug/libfeedbug.so
First feed_input + run: OK
Traceback (most recent call last):
  File "/mnt/tdm-dic/users/cesco-safe/code/feed-bug/main.py", line 85, in main
    pipe.feed_input("Reader", batch2)  # ← crash here
    ~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^
  File "/mnt/tdm-dic/users/cesco-safe/code/feed-bug/.venv/lib/python3.13/site-packages/nvidia/dali/pipeline.py", line 1337, in feed_input
    self._feed_input(name, data, layout, cuda_stream, use_copy_kernel)
    ~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/mnt/tdm-dic/users/cesco-safe/code/feed-bug/.venv/lib/python3.13/site-packages/nvidia/dali/pipeline.py", line 1248, in _feed_input
    self._pipe.SetExternalTLInput(name, data, cuda_stream_ptr, use_copy_kernel)
    ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: CUDA runtime API error cudaErrorInvalidResourceHandle (400):
invalid resource handle

The above exception was the direct cause of the following exception:

Traceback (most recent call last):
  File "/mnt/tdm-dic/users/cesco-safe/code/feed-bug/main.py", line 99, in <module>
    main()
    ~~~~^^
  File "/mnt/tdm-dic/users/cesco-safe/code/feed-bug/main.py", line 90, in main
    raise RuntimeError(
    ...<4 lines>...
    ) from e
RuntimeError: BUG REPRODUCED: Second feed_input crashed.
This error is caused by DALI's AccessOrder::wait() when cudaGetDevice() is called in the host-order branch.

Original error: CUDA runtime API error cudaErrorInvalidResourceHandle (400):
invalid resource handle

Other/Misc.

No response

Check for duplicates

• I have searched the open bugs/issues and have found no duplicates for this bug report

[JanuszL]

Hi @fversaci,

Thanks for the reproducer. We were able to reproduce the crash and this looks like an ABI incompatibility between the DALI build and the external plugin.

Short technical summary: the plugin derives from DALI's C++ InputOperator. That class embeds DALI implementation types whose layout depends on private DALI build configuration. In particular, NVML_ENABLED changes the layout of WorkerThread, which shifts later InputOperator members such as internal_copy_order_.

If DALI itself is built with NVML_ENABLED=1, but the external plugin is built without that define, the plugin constructor initializes internal_copy_order_ at one offset while DALI runtime code reads it from another offset. The value read by DALI is then effectively uninitialized/corrupted and can be interpreted as a CUDA stream, which leads to the crash during the second feed_input when
DALI recycles the input buffer and updates its access order.

We have an idea how to address this on the DALI side by hiding the affected thread-pool implementation details from public headers so private build flags do not change the layout seen by external plugins but we need to discuss pross and cons. The proposed draft is #6362

[fversaci]

Hi @JanuszL

thanks for looking into this so quickly! I'll follow the proposed draft.

[JanuszL]

It has been merged and is available in the nightly build. Please check if it works as you expect.