[CI] New test failures on GB300

[jcrist]

We're seeing some new test failures on the newly rolled out GB300 architecture in CI.

Some failures are test failures (incorrect result/result outside the previously accepted bounds). These may be bugs, or just tests that need a bit of a tolerance nudge.

Some test failures are full on crashes. These are almost certainly bugs, but may not be specific to cuml.

The failures are reproducible, here's a link to the latest nightly run: https://github.com/NVIDIA/cuml/actions/runs/32707368663/job/97371258132

[jcrist]

(Assigned @betatim just for initial diagnosis/triage. Thanks!)

[betatim]

Some of the HDBSCAN failures should be fixed by NVIDIA/cuvs#2503. The code was setting counters to random values, then adding to them a bunch of times and asserting that their value equaled the number of additions. Because they started at random values instead of zero that assert fails.

These two tests cuml_accel_tests/test_basic_estimators.py::test_hdbscan and cuml_accel_tests/integration/test_sklearn_hdbscan.py::test_hdbscan_copy_default_warning won't be fixed by this, I think.

Both run HDBSCAN().fit(X) with default parameters on make_blobs(n_samples=100, centers=3, random_state=42), and both fail inside lib.hdbscan() with:

RuntimeError: copy_if failed on 2nd step: cudaErrorIllegalAddress: an illegal memory access was encountered

I think this is a real illegal address, not a side effect of some other test. With #8519 we will know for sure. I don't think this is the same kind of problem as the above PRs that are all related to the change in RMM and us initialising things incorrectly. I asked AI to investigate this and it concluded "I can't work it out". So, if you'd like to show off that you are smarter than AI this is a good thing to fix :D It did give me this:

Does not reproduce on x86. Both tests pass in 0.3s on an RTX A6000 with CUDA 13.3.

The remaining candidates are the three thrust::copy_if calls in the HDBSCAN code: hdbscan/detail/utils.h:115, hdbscan/condensed_hierarchy.cu:164 and hdbscan/detail/select.cuh:320. The first looks most suspicious, since its output buffers are sized by a separate transform_reduce over the same a > 1 predicate, so any disagreement between the count and the copy overruns them. Next step is compute-sanitizer on a GB300.

I don't have a GB300, so this is the end of the road of chasing nightly failures for me today. Someone who knows HDBSCAN should take over.

[betatim]

Dante will take a look at the HDBSCAN + GB300 issue on Wednesday.

[dantegd]

I reproduced this on GB300 with CUDA 13.3, driver 595.58.03, eight pytest-xdist workers using --dist worksteal, and CUDA_LAUNCH_BLOCKING unset.

The copy_if failed on 2nd step: cudaErrorIllegalAddress exception is a downstream symptom rather than a bug in that copy_if. I found two independent cross-stream ordering bugs in the HDBSCAN dependency path:

1. RAFT MST weight alteration creates a cuRAND generator without assigning it to the MST stream. Random values can therefore be generated on stream 0 while the alteration kernel consumes them on the handle stream. Fix is in PR Fix curand stream ordering in MST alteration raft#3125 in RAFT.

2. In cuVS brute-force KNN, fusedL2Knn can run on a selected stream-pool stream, while the L2SqrtExpandedpost-processing map is launched using the main handle. There is no dependency making the main stream wait for the selected stream before reading the KNN output. Fix is in Fix stream ordering for brute-force L2 sqrt post-processing cuvs#2509 in cuVS

The local test results were:

RAFT fix	cuVS fix	Result
No	No	Reproduces/Fails
Yes	No	Reproduces/Fails
No	Yes	Reproduces/Fails
Yes	Yes	Passes

With both fixes and launch blocking still disabled:

• Eight processes × 100 HDBSCAN fits: 800/800 passed.
• cuml_accel_tests/test_basic_estimators.py::test_hdbscan: passed.
• cuml_accel_tests/integration/test_sklearn_hdbscan.py::test_hdbscan_copy_default_warning: passed.
• The full accel workload completed without either HDBSCAN failure.

[betatim]

PRs and issues to track are the ones that refer to this issue as well as

NVIDIA/cuvs#2504

[jcrist]

This is now resolved. Thanks everyone!