Fix RTX PRO 6000 Blackwell CI

[bdice]

Description

This reverts commit 9f31e1d. Closes #21953.

With rapidsai/rapids-cmake#996, the previous bug should be fixed.

Checklist

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

[davidwendt]

I'm inclined to believe the new rtxpro6000 test failure is due to a compute-sanitizer bug or incompatibility.

[davidwendt]

Verified this is likely a compute-sanitizer issue. The error appears to be a false-positive on the TMA-based operation cp.async.bulk. This may be fixed in a newer compute-sanitizer (2026.1) which I think may be part of 13.2.

We could modify https://github.com/rapidsai/cudf/blob/main/ci/run_cudf_examples.sh to run the examples without compute-sanitizer when running on rtxpro6000 but I'm not sure how to detect that.

[davidwendt]

Installed CUDA toolkit 13.2 and compute-sanitizer 2026.1 and the error still persists.
More detailed error for the hybrid_scan_io example:

========= COMPUTE-SANITIZER
========= Invalid __shared__ read of size 16 bytes
=========     at void cuda::ptx::__4::cp_async_bulk_cp_mask<void>(cuda::std::__4::integral_constant<cuda::ptx::__4::dot_space, (cuda::ptx::__4::dot_space)0>, cuda::std::__4::integral_constant<cuda::ptx::__4::dot_space, (cuda::ptx::__4::dot_space)2>, void *, const void *, const unsigned int &, const unsigned short &)+0x5400 in cp_async_bulk.h:236
=========     by thread (128,0,0) in block (0,0,0)
=========     Access to 0x400 is out of bounds
=========         Device Frame: void cub::detail::warpspeed::squadStoreBulkSync<int>(cub::detail::warpspeed::Squad, cub::detail::warpspeed::CpAsyncOobInfo<T1>, const cuda::std::byte *)+0x52e0 in load_store.cuh:288
=========         Device Frame: void cub::detail::scan::kernelBody<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, int, int, int, cuda::std::__4::plus<void>, int, (bool)0>(cub::detail::warpspeed::Squad, cub::detail::warpspeed::SpecialRegisters, const cub::detail::scan::scanKernelParams<T2, T3, T4> &, T5, T6)+0x4fc0 in kernel_scan_warpspeed.cuh:721
=========         Device Frame: void cub::detail::scan::device_scan_lookahead_body<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, (bool)0, int, int, int, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>>(cub::detail::scan::scanKernelParams<T4, T5, T6>, T7, const T8 &)::[lambda(cub::detail::warpspeed::Squad) (instance 1)]::operator ()(cub::detail::warpspeed::Squad) const+0x3b80 in kernel_scan_warpspeed.cuh:810
=========         Device Frame: void cub::detail::warpspeed::squadDispatch<(int)1, void cub::detail::scan::device_scan_lookahead_body<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, (bool)0, int, int, int, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>>(cub::detail::scan::scanKernelParams<T4, T5, T6>, T7, const T8 &)::[lambda(cub::detail::warpspeed::Squad) (instance 1)]>(cub::detail::warpspeed::SpecialRegisters, const cub::detail::warpspeed::SquadDesc (&)[T1], T2, int)+0x3b70 in squad.cuh:114
=========         Device Frame: void cub::detail::warpspeed::squadDispatch<(int)2, void cub::detail::scan::device_scan_lookahead_body<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, (bool)0, int, int, int, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>>(cub::detail::scan::scanKernelParams<T4, T5, T6>, T7, const T8 &)::[lambda(cub::detail::warpspeed::Squad) (instance 1)]>(cub::detail::warpspeed::SpecialRegisters, const cub::detail::warpspeed::SquadDesc (&)[T1], T2, int)+0x3b70 in squad.cuh:145
=========         Device Frame: void cub::detail::warpspeed::squadDispatch<(int)2, void cub::detail::scan::device_scan_lookahead_body<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, (bool)0, int, int, int, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>>(cub::detail::scan::scanKernelParams<T4, T5, T6>, T7, const T8 &)::[lambda(cub::detail::warpspeed::Squad) (instance 1)]>(cub::detail::warpspeed::SpecialRegisters, const cub::detail::warpspeed::SquadDesc (&)[T1], T2, int)+0xd0 in squad.cuh:135
=========         Device Frame: void cub::detail::warpspeed::squadDispatch<(int)5, void cub::detail::scan::device_scan_lookahead_body<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, (bool)0, int, int, int, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>>(cub::detail::scan::scanKernelParams<T4, T5, T6>, T7, const T8 &)::[lambda(cub::detail::warpspeed::Squad) (instance 1)]>(cub::detail::warpspeed::SpecialRegisters, const cub::detail::warpspeed::SquadDesc (&)[T1], T2, int)+0xb0 in squad.cuh:135
=========         Device Frame: void cub::detail::scan::device_scan_lookahead_body<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, (bool)0, int, int, int, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>>(cub::detail::scan::scanKernelParams<T4, T5, T6>, T7, const T8 &)+0x90 in kernel_scan_warpspeed.cuh:809
=========         Device Frame: void cub::detail::scan::DeviceScanKernel<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void>>, int *, int *, cub::ScanTileState<int, (bool)1>, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int *>, unsigned int, int, (bool)0, int>(T2, T3, cub::detail::scan::tile_state_kernel_arg_t<T4, T8>, int, T5, T6, T7, int)+0x10 in kernel_scan.cuh:207
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cuLaunchKernel [0x39d6c4] in libcuda.so.1
=========         Host Frame:  [0x14cfc] in libcudart.so.13
=========         Host Frame: cudaLaunchKernel [0x7fea7] in libcudart.so.13
=========         Host Frame: void cub::detail::scan::DeviceScanKernel<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void> >, int*, int*, cub::ScanTileState<int, true>, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int*>, unsigned int, int, false, int>(int*, int*, cub::detail::scan::tile_state_kernel_arg_t<cub::ScanTileState<int, true>, int>, int, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int*>, unsigned int, int) [0x1167136] in libcudf.so
=========         Host Frame: auto cub::detail::scan::dispatch<(cub::ForceInclusive)1, int*, int*, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int*>, unsigned int, int, cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void> >, cub::detail::scan::DeviceScanKernelSource<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void> >, int*, int*, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int*>, unsigned int, int, (cub::ForceInclusive)1>, cub::detail::TripleChevronFactory>(void*, unsigned long&, int*, int*, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int*>, unsigned int, CUstream_st*, cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void> >, cub::detail::scan::DeviceScanKernelSource<cub::detail::scan::policy_selector_from_types<int, int, int, unsigned int, cuda::std::__4::plus<void> >, int*, int*, cuda::std::__4::plus<void>, cub::detail::InputValue<int, int*>, unsigned int, int, (cub::ForceInclusive)1>, cub::detail::TripleChevronFactory)::{lambda(auto:1)#1}::operator()<cuda::std::__4::integral_constant<cub::detail::scan::scan_policy, cub::detail::scan::scan_policy{384, 22, (cub::BlockLoadAlgorithm)4, (cub::CacheLoadModifier)0, (cub::BlockStoreAlgorithm)4, (cub::BlockScanAlgorithm)2, cub::detail::delay_constructor_policy{(cub::detail::delay_constructor_kind)6, 1904u, 830u}, cub::detail::scan::scan_warpspeed_policy{true, 4, 4, 1, 1, 1, 4, 352, 63, 8064}}> >(cuda::std::__4::integral_constant<cub::detail::scan::scan_policy, cub::detail::scan::scan_policy{384, 22, (cub::BlockLoadAlgorithm)4, (cub::CacheLoadModifier)0, (cub::BlockStoreAlgorithm)4, (cub::BlockScanAlgorithm)2, cub::detail::delay_constructor_policy{(cub::detail::delay_constructor_kind)6, 1904u, 830u}, cub::detail::scan::scan_warpspeed_policy{true, 4, 4, 1, 1, 1, 4, 352, 63, 8064}}>) const [clone .isra.0] [0x1169297] in libcudf.so
=========         Host Frame: cudf::io::parquet::detail::decode_page_headers(cudf::io::parquet::detail::pass_intermediate_data&, cuda::std::__4::span<cudf::io::parquet::detail::PageInfo, 18446744073709551615ul>, bool, rmm::cuda_stream_view) [0x1174b66] in libcudf.so
=========         Host Frame: cudf::io::parquet::experimental::detail::hybrid_scan_reader_impl::setup_compressed_data(cudf::host_span<cuda::std::__4::span<unsigned char const, 18446744073709551615ul> const, 18446744073709551615ul>) [0x1036614] in libcudf.so
=========         Host Frame: cudf::io::parquet::experimental::detail::hybrid_scan_reader_impl::setup_next_pass(cudf::host_span<cuda::std::__4::span<unsigned char const, 18446744073709551615ul> const, 18446744073709551615ul>) [0x1012ce6] in libcudf.so
=========         Host Frame: cudf::io::parquet::experimental::detail::hybrid_scan_reader_impl::handle_chunking(cudf::io::parquet::detail::reader_impl::read_mode, cudf::host_span<cuda::std::__4::span<unsigned char const, 18446744073709551615ul> const, 18446744073709551615ul>, cudf::host_span<bool const, 18446744073709551615ul>) [0x10151aa] in libcudf.so
=========         Host Frame: cudf::io::parquet::experimental::detail::hybrid_scan_reader_impl::prepare_data(cudf::io::parquet::detail::reader_impl::read_mode, cudf::host_span<std::vector<int, std::allocator<int> > const, 18446744073709551615ul>, cudf::host_span<cuda::std::__4::span<unsigned char const, 18446744073709551615ul> const, 18446744073709551615ul>, cudf::host_span<bool const, 18446744073709551615ul>) [0x10257f5] in libcudf.so
=========         Host Frame: cudf::io::parquet::experimental::detail::hybrid_scan_reader_impl::materialize_filter_columns(cudf::host_span<std::vector<int, std::allocator<int> > const, 18446744073709551615ul>, cudf::host_span<cuda::std::__4::span<unsigned char const, 18446744073709551615ul> const, 18446744073709551615ul>, cudf::mutable_column_view&, cudf::io::parquet::experimental::use_data_page_mask, cudf::io::parquet_reader_options const&, rmm::cuda_stream_view, rmm::detail::cccl_async_resource_ref<cuda::mr::__4::__version_bump_ver4_::resource_ref<cuda::mr::__4::device_accessible> >) [0x1032ff5] in libcudf.so
=========         Host Frame: cudf::io::parquet::experimental::hybrid_scan_reader::materialize_filter_columns(cudf::host_span<int const, 18446744073709551615ul>, cudf::host_span<cuda::std::__4::span<unsigned char const, 18446744073709551615ul> const, 18446744073709551615ul>, cudf::mutable_column_view&, cudf::io::parquet::experimental::use_data_page_mask, cudf::io::parquet_reader_options const&, rmm::cuda_stream_view, rmm::detail::cccl_async_resource_ref<cuda::mr::__4::__version_bump_ver4_::resource_ref<cuda::mr::__4::device_accessible> >) const [0x100e4a9] in libcudf.so
=========         Host Frame: std::unique_ptr<cudf::table, std::default_delete<cudf::table> > hybrid_scan<false, true>(io_source const&, std::optional<cudf::ast::operation const>, std::unordered_set<hybrid_scan_filter_type, std::hash<hybrid_scan_filter_type>, std::equal_to<hybrid_scan_filter_type>, std::allocator<hybrid_scan_filter_type> > const&, bool, rmm::cuda_stream_view, rmm::detail::cccl_async_resource_ref<cuda::mr::__4::__version_bump_ver4_::resource_ref<cuda::mr::__4::device_accessible> >) [0x3663b] in hybrid_scan_io
=========         Host Frame: main [0xe893] in hybrid_scan_io
=========

This output includes line numbers too.

[bdice]

CI output is now more useful with 4254f63

Running basic example: compute-sanitizer --tool memcheck ./basic_example

[bdice]

/merge