Sync the stream before a GPU video decoder frees its frame buffers by jantonguirao · Pull Request #6556 · NVIDIA/DALI

jantonguirao · 2026-10-27T15:37:47Z

Category:

Bug fix (non-breaking change which fixes an issue)

Description:

This PR is split out of #6545, which was too large to review as one change. It is based directly on main and does not depend on the other PRs of the split.

Problem. FramesDecoderGpu keeps its decoded frames in frame_buffer_ (DeviceBuffers), and copies frames out of it into the output with asynchronous device-to-device copies (copyD2D) enqueued on the decoder's stream. When the decoder is destroyed, frame_buffer_ is returned to the memory pool with a plain (not stream-ordered) deallocation, while one of those copies may still be pending on the stream. The memory can then be handed out to work on another stream and overwritten before the pending copy has read it, so the copied frame may contain wrong data. Nothing reports this: the result is silently corrupted output, not an error.

When it happens. fn.experimental.readers.video (GPU) replaces its decoder in Prefetch every time the next sample comes from another file, so a decoder is destroyed right after the last copy out of it has been enqueued, which makes the window real. The same destructor runs for every other user of FramesDecoderGpu: fn.experimental.decoders.video and fn.experimental.inputs.video (mixed backend). The decode lanes of #6545 (#6563, #6564) run several GPU decoders concurrently, which makes the race more likely, but it exists on main already.

Fix. The destructor of FramesDecoderGpu now calls cudaStreamSynchronize(stream_) before its members (including frame_buffer_) are released, so every copy out of the frame buffers has finished when their memory goes back to the pool. The call is checked with CUDA_DTOR_CALL, as other CUDA calls in DALI destructors: only the error of an unloading CUDA runtime is ignored.

Additional information:

Affected modules and functionalities:

dali/operators/video/frames_decoder_gpu.cc: FramesDecoderGpu::~FramesDecoderGpu() synchronizes stream_ before the frame buffers are freed.

Key points relevant for the review:

The cost is one stream synchronization per decoder destruction (e.g. one per file switch in the reader), not per frame. It waits for all work enqueued on the decoder's stream so far (the stream the operator passed to the decoder), not only for the copies out of frame_buffer_.
cudaStreamSynchronize in the destructor is checked with CUDA_DTOR_CALL. Apart from an unloading runtime, it can only fail with a sticky error from earlier work on the stream, when the CUDA context is unusable anyway.

Tests:

Existing tests apply
- frames_decoder_test.cc: FramesDecoderGpuTest.* (decoders created and destroyed on a stream)
- video_reader_decoder_op_test.cc: VideoReaderDecoderGpuTest.*, VideoReaderDecoderCompareTest.CompareReaders
- test_video_reader.py: test_compare_experimental_to_legacy_reader_*, test_uniform_sample with device="gpu" (several files per reader, so decoders are destroyed while reading)
- decoder/test_video.py: fn.experimental.decoders.video tests with device="mixed"
- input/test_video.py: test_video_input_* with device="mixed"
New tests added
- Python tests
- GTests
- Benchmark
- Other
N/A

Checklist

Documentation

DALI team only

Requirements

Implements new requirements
Affects existing requirements
N/A

REQ IDs: N/A

JIRA TASK: DALI-4935

FramesDecoderGpu's frame_buffer_ is returned to the memory pool without stream ordering when the decoder is destroyed (the reader replaces its decoder on every file switch), while an asynchronous copy out of it may still be pending on the decoder's stream. Another stream could then be given that memory. The destructor now synchronizes the stream first. Signed-off-by: Joaquin Anton Guirao <janton@nvidia.com>

copy-pr-bot · 2026-10-27T15:37:51Z

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.

Copilot

🟢 Approval recommended

The synchronization precedes frame-buffer destruction and addresses the identified lifetime race, with no blocking issues found in static review.

No open issues 🎉

What changed in this PR

This focused fix prevents GPU video frame buffers from returning to the memory pool while asynchronous copies still use them. It addresses the existing decoder lifetime race independently of the parallel decoding changes.

Changes:

Synchronize the decoder’s CUDA stream before releasing frame buffers.
Clear synchronization errors without throwing from the destructor.

File	Description
`dali/operators/video/frames_decoder_gpu.cc`	Wait for pending stream work before decoder members are destroyed.

_{Review effort: Balanced}

greptile-apps · 2026-10-27T15:42:30Z

Confidence Score: 4/5

[Medium impact] The change reaches GPU video decoding workflows that use this decoder.

Restore non-throwing cleanup before merging so a CUDA error cannot abort the application during decoder destruction.

Findings

GPU errors abort the application ▶

Summary

The PR waits for the decoder's stream before returning its frame buffers to the memory pool.

GPU video decoders finish queued copies before releasing frame buffers.

_{Reviews (2) · Last reviewed commit: "Check the stream sync in the GPU video d..." · Reviewed by Greptile}

rostan-t · 2026-10-27T15:46:26Z

+  if (cudaStreamSynchronize(stream_) != cudaSuccess)
+    (void)cudaGetLastError();


Nitpick: shouldn't this be CUDA_DTOR_CALL?

Yes, done in aec9d1a.

mzient · 2026-10-27T15:52:08Z

+  // for the copies out of it (CopyFrame, copyD2D) that may still be pending on stream_. Nothing
+  // is thrown from the destructor: a failing sync means the CUDA context is already broken, and
+  // that error is reported by the call that caused it.
+  if (cudaStreamSynchronize(stream_) != cudaSuccess)


What kind of error is acceptable here, except the "unloading" one already swallowed by CUDA_DTOR_CALL?

None that should be hidden. Apart from unloading, the sync can only fail here with a sticky error from earlier work on the stream, and then the context is unusable anyway. I had swallowed it to avoid throwing from a destructor, but DALI checks CUDA calls in destructors with CUDA_DTOR_CALL elsewhere (e.g. cudaStreamDestroy and cudaEventDestroy in CUDAStream/CUDAEvent), so I switched to it in aec9d1a.

…_DTOR_CALL As the other destructors that synchronize a stream do: only the error of an unloading CUDA runtime is ignored. Signed-off-by: Joaquin Anton Guirao <janton@nvidia.com>

jantonguirao · 2026-10-27T15:57:15Z

!build

greptile-apps · 2026-10-27T16:00:37Z

 FramesDecoderGpu::~FramesDecoderGpu() {
+  // frame_buffer_ goes back to the memory pool right after this, without stream ordering: wait
+  // for the copies out of it (CopyFrame, copyD2D) that may still be pending on stream_
+  CUDA_DTOR_CALL(cudaStreamSynchronize(stream_));


GPU errors abort the application

[Bug] CUDA_DTOR_CALL throws for every CUDA failure except runtime unloading. If cudaStreamSynchronize reports an earlier GPU failure during decoder cleanup, that exception leaves the non-throwing FramesDecoderGpu destructor and calls std::terminate. The host process aborts instead of reporting the pipeline error.

Restore the previous non-throwing check so cleanup clears the CUDA error without throwing.

Suggested change

CUDA_DTOR_CALL(cudaStreamSynchronize(stream_));

if (cudaStreamSynchronize(stream_) != cudaSuccess)

(void)cudaGetLastError();

We are OK with that

dali-automaton · 2026-10-27T16:02:43Z

CI MESSAGE: [77044862]: BUILD STARTED

dali-automaton · 2026-10-27T18:48:26Z

CI MESSAGE: [77044862]: BUILD PASSED

Copilot AI balanced review requested due to automatic review settings October 27, 2026 15:37

Copilot started reviewing on behalf of jantonguirao October 27, 2026 15:38

View session

This was referenced Oct 27, 2026

Do not wait on the host for every picture the GPU video decoder converts #6561

Draft

Do not map, convert nor decode video pictures that are never read #6562

Draft

jantonguirao added the trivial change label Oct 27, 2026

This was referenced Oct 27, 2026

Decode the video reader's batches with up to num_threads decode lanes #6563

Draft

Run the video reader's decode lanes on their own threads, a batch ahead #6564

Draft

Copilot AI reviewed Oct 27, 2026

View reviewed changes

rostan-t approved these changes Oct 27, 2026

View reviewed changes

mzient reviewed Oct 27, 2026

View reviewed changes

Check the stream sync in the GPU video decoder's destructor with CUDA…

aec9d1a

…_DTOR_CALL As the other destructors that synchronize a stream do: only the error of an unloading CUDA runtime is ignored. Signed-off-by: Joaquin Anton Guirao <janton@nvidia.com>

greptile-apps Bot reviewed Oct 27, 2026

View reviewed changes

jantonguirao merged commit 78e45ee into NVIDIA:main Oct 27, 2026
10 checks passed

		if (cudaStreamSynchronize(stream_) != cudaSuccess)
		(void)cudaGetLastError();

	CUDA_DTOR_CALL(cudaStreamSynchronize(stream_));
	if (cudaStreamSynchronize(stream_) != cudaSuccess)
	(void)cudaGetLastError();

Conversation

jantonguirao commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Category:

Description:

Additional information:

Affected modules and functionalities:

Key points relevant for the review:

Tests:

Checklist

Documentation

DALI team only

Requirements

Uh oh!

copy-pr-bot Bot commented Oct 27, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟢 Approval recommended

Uh oh!

greptile-apps Bot commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Confidence Score: 4/5

Findings

Summary

Uh oh!

rostan-t Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

jantonguirao Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

mzient Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jantonguirao Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jantonguirao commented Oct 27, 2026

Uh oh!

greptile-apps Bot Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

jantonguirao Oct 27, 2026

Choose a reason for hiding this comment

Uh oh!

dali-automaton commented Oct 27, 2026

Uh oh!

dali-automaton commented Oct 27, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

jantonguirao commented Oct 27, 2026 •

edited

Loading

greptile-apps Bot commented Oct 27, 2026 •

edited

Loading

mzient Oct 27, 2026 •

edited

Loading

jantonguirao Oct 27, 2026 •

edited

Loading