Fix GB300 segfaults from destroying an async memory pool with pending frees in test_arrow_object_lifetime by bdice · Pull Request #24688 · NVIDIA/cudf

bdice · 2026-10-19T15:16:50Z

Description

Fixes the GB300-only segfaults in pylibcudf's parquet tests (e.g. test_read_parquet_basic, nightly example). These are not caused by #24590.

Root cause

test_interop.py::test_arrow_object_lifetime installs a temporary rmm.mr.CudaAsyncMemoryResource, allocates a table from it, restores the previous resource, and lets the async resource, and therefore its stream-ordered memory pool, be destroyed while GPU work and frees on that pool are still pending. After that, the next nvCOMP 6 hardware decompression (the Blackwell Decompression Engine, which only GB300 has in our CI) in the same process segfaults inside the CUDA driver at a later cuStreamSynchronize. It started with the nvCOMP 6.0 upgrade (#24580). It needs other processes using the GPU concurrently, so it showed up as a random parquet test crashing in whichever pytest-xdist worker had run this test earlier.

How it was isolated (all on GB300 in this PR's CI):

Forcing the CUDA decompression backend removes the crash, and the same build with the hardware backend crashes. Legacy vs. per-thread default stream, kvikio threading, nvCOMP's pinned pool, and dependency versions are not factors.
Delta-debugging a crashing worker's exact 2,311-test history, replayed in 8 concurrent processes, reduced it to test_arrow_object_lifetime followed by a parquet read.
ci/gb300_repro.py reproduces it without cuDF's test suite, and its lowest layer uses only cuda-python plus a pylibcudf parquet read (8 concurrent processes x 3 rounds):

setup before the parquet read	processes that segfault
none	0/24
RMM async MR + table, destroyed with frees pending	15–19/24
same, `cudaDeviceSynchronize()` before destroy	0/24
raw cuda-python pool: async writes + frees, destroyed with work pending	24/24
same, synchronized before destroy	0/24
same, pending, CUDA backend instead of hardware decompression	0/24

The underlying bug is in the CUDA driver / nvCOMP 6 hardware decompression path. Destroying a pool with outstanding work is allowed by the CUDA API, since its release is deferred.

Fix

Synchronize the device before the test's temporary async resource is released, so its pool is never destroyed with work pending. The test now also asserts directly, via a weak reference, that the table keeps the resource alive, instead of relying on a segfault to detect a lifetime bug. No library changes are needed, and hardware decompression stays enabled.

Validation (GB300, full pylibcudf suite with `--numprocesses=8`)

A/B where only test_interop.py differs: original 2/6 runs crash, fixed 0/10 (11,585 passed each run).
gb300-validate (6 full-suite runs per head): 0 crashes or failures on every head validated (24 runs total), and the GB300 wheel-tests-cudf job passes.

Follow-ups outside this PR: report the driver/nvCOMP bug, and consider having RMM's cuda_async_memory_resource synchronize before cudaMemPoolDestroy.

Note on CI configuration

As requested for this investigation, this PR's .github/workflows/pr.yaml runs tests only on GB300: an arm64 CUDA 13 wheel build on CPU runners, the pylibcudf wheel tests, and ci/gb300_validate.sh (6 full-suite runs that fail on any crash or test failure, plus the ci/gb300_repro.py reproduction matrix). Revert pr.yaml and delete ci/gb300_* before merging; the fix itself is only python/pylibcudf/tests/test_interop.py.

Checklist

I am familiar with the Contributing Guidelines.
New or existing tests cover these changes.
The documentation is up to date with these changes.

…eels

copy-pr-bot · 2026-10-19T15:16:54Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

bdice · 2026-10-19T15:16:57Z

/ok to test 1764999

bdice · 2026-10-19T16:28:15Z

/ok to test ccb8065

…toggles

bdice · 2026-10-19T16:54:36Z

/ok to test ba8515b

…ecompression

bdice · 2026-10-19T17:18:42Z

/ok to test 7baf366

…state

bdice · 2026-10-19T17:25:00Z

/ok to test e392838

bdice · 2026-10-19T17:31:10Z

/ok to test d7c7601

bdice · 2026-10-19T17:42:12Z

/ok to test 883c4f6

…pytest shards

bdice · 2026-10-19T17:49:57Z

/ok to test dac3b03

bdice · 2026-10-19T17:58:09Z

/ok to test 77441db

bdice · 2026-10-19T18:04:56Z

/ok to test 3f56bab

bdice · 2026-10-19T18:12:51Z

/ok to test 24c9e7e

bdice · 2026-10-19T18:48:25Z

/ok to test 24c9e7e

bdice · 2026-10-19T19:02:40Z

/ok to test 536941f

… and nvCOMP pinned pool

bdice · 2026-10-19T22:25:49Z

/ok to test c55ab88

… a DE read

bdice · 2026-10-19T22:37:16Z

/ok to test ed82f19

bdice · 2026-10-19T22:44:16Z

/ok to test b02e3cf

…object_lifetime

bdice · 2026-10-19T22:52:08Z

/ok to test aa7fad7

bdice · 2026-10-19T23:09:46Z

/ok to test 9c09137

bdice · 2026-10-19T23:25:23Z

/ok to test a05f967

bdice · 2026-10-19T23:40:56Z

/ok to test 81c3a6c

bdice · 2026-10-19T23:56:25Z

/ok to test 81c3a6c

bdice · 2026-10-19T23:56:51Z

/ok to test 60d9966

bdice · 2026-10-20T00:12:58Z

/ok to test 17ad09a

bdice · 2026-10-20T00:28:22Z

/ok to test acaf7cf

Bisect test_read_parquet_basic segfault on GB300 using main branch wh…

1764999

…eels

bdice added the DO NOT MERGE Hold off on merging; see PR for details label Oct 19, 2026

github-actions Bot assigned bdice Oct 19, 2026

Bisect GB300 segfault across rmm, kvikio, and cuda-core versions

ccb8065

Probe GB300 segfault with decompression, kvikio, and launch-blocking …

ba8515b

…toggles

Add standalone GB300 parquet repro for path vs bytes sources and HW d…

7baf366

…ecompression

Run pylibcudf tests serially on GB300 to find the test that corrupts …

e392838

…state

Capture the test sequence of the crashing xdist worker on GB300

d7c7601

Bisect which preceding pylibcudf test module triggers the GB300 segfault

883c4f6

Capture all-thread backtraces of the GB300 segfault under concurrent …

dac3b03

…pytest shards

Run GB300 xdist workers under gdb to capture all-thread backtraces

77441db

Capture GB300 segfault address, registers, and crashing thread first

3f56bab

Retry GB300 xdist-under-gdb up to 10 times

24c9e7e

Make nvCOMP hardware decompression opt-in to fix GB300 segfaults

536941f

github-actions Bot added the libcudf Affects libcudf (C++/CUDA) code. label Oct 19, 2026

Revert opt-in HW decompression; probe DE crash with per-thread stream…

b5b3daa

… and nvCOMP pinned pool

Bound ddmin probe time and log reduction progress

c55ab88

Repro GB300 segfault by destroying a HW-decompress memory pool before…

ed82f19

… a DE read

Narrow the GB300 repro to deferred destruction of a HW-decompress pool

b02e3cf

Synchronize before releasing the async memory resource in test_arrow_…

aa7fad7

…object_lifetime

github-actions Bot added Python Affects Python cuDF API. pylibcudf Issues specific to the pylibcudf package labels Oct 19, 2026

github-project-automation Bot added this to cuDF Python Oct 19, 2026

github-project-automation Bot moved this to Todo in cuDF Python Oct 19, 2026

Use a weakref-able async resource subclass in test_arrow_object_lifetime

9c09137

Remove GB300-only CI matrix and debugging scripts

58833ae

bdice removed the DO NOT MERGE Hold off on merging; see PR for details label Oct 19, 2026

bdice changed the title ~~Fix test_read_parquet_basic segfault on GB300~~ Fix GB300 segfaults from destroying an async memory pool with pending frees in test_arrow_object_lifetime Oct 19, 2026

bdice added bug Something isn't working non-breaking Non-breaking change labels Oct 19, 2026

Restore GB300-only CI matrix with fix validation

a05f967

Add standalone reproduction of the GB300 memory pool and DE crash

81c3a6c

Extend GB300 repro to written raw pools and RMM-only buffers

60d9966

Print failing GB300 repro process output

17ad09a

Add sync and DE-off controls to the raw pool repro

acaf7cf

Conversation

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

Uh oh!

Description

Root cause

Fix

Validation (GB300, full pylibcudf suite with --numprocesses=8)

Note on CI configuration

Checklist

Uh oh!

copy-pr-bot Bot commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 19, 2026

Uh oh!

bdice commented Oct 20, 2026

Uh oh!

bdice commented Oct 20, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

bdice commented Oct 19, 2026 •

edited

Loading

Validation (GB300, full pylibcudf suite with `--numprocesses=8`)