Support CUDA 12.2

[jameslamb]

Description

• switches to CUDA 12.2.2 for building wheels
• adds new tests running against CUDA 12.2.2
• adds some "12.*" globs in dependencis.yaml missed in refactor CUDA versions in dependencies.yaml #5712

Closes #5765

Notes for Reviewers

This is part of ongoing work to build and test wheels against CUDA 12.2.2 across all of RAPIDS. For more details see:

• Add support for CUDA 12.2 wheels build-planning#7
• add CUDA 12.2 conda and wheel jobs shared-workflows#166

(created with rapids-reviser)

[jameslamb]

The new CUDA 12.2 builds are failing with 2 issues.

conda packages

add_library cannot create imported target "CCCL::Thrust" because another
  target with the same name already exists.

(build link)

That should be fixed by rapidsai/rapids-cmake#522.

wheels

  FAILED: cuml-cpp/CMakeFiles/cuml++.dir/src/explainer/tree_shap.cu.o
  sccache /usr/local/cuda/bin/nvcc -forward-unknown-to-host-compiler -DCUML_CPP_API -DCUML_ENABLE_GPU -DCUTLASS_ENABLE_CUDNN=1 -DCUTLASS_NAMESPACE=raft_cutlass -DDISABLE_CUSPARSE_DEPRECATED -DFMT_HEADER_ONLY=1 -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE -DRAFT_COMPILED -DRAFT_SYSTEM_LITTLE_ENDIAN=1 -DRMM_STATIC_CUDART -DSPDLOG_FMT_EXTERNAL -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -Dcuml___EXPORTS -I/__w/cuml/cuml/cpp/include -I/__w/cuml/cuml/cpp/src -I/__w/cuml/cuml/cpp/src/metrics -I/__w/cuml/cuml/cpp/src_prims -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/raft-src/cpp/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/treelite-src/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/treelite-build/include -I/__w/cuml/cuml/cumlprims_mg/cpp/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/rmm-src/include -I/usr/local/cuda/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/fmt-src/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/spdlog-src/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/cuco-src/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/nvidiacutlass-src/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/nvidiacutlass-build/include -I/__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/gputreeshap-src -isystem /usr/local/cuda/targets/sbsa-linux/include -O3 -DNDEBUG -std=c++17 "--generate-code=arch=compute_70,code=[sm_70]" "--generate-code=arch=compute_75,code=[sm_75]" "--generate-code=arch=compute_80,code=[sm_80]" "--generate-code=arch=compute_86,code=[sm_86]" "--generate-code=arch=compute_90,code=[compute_90,sm_90]" -Xcompiler=-fPIC --expt-extended-lambda --expt-relaxed-constexpr -Werror=all-warnings -Xcompiler=-Wall,-Werror,-Wno-error=deprecated-declarations,-Wno-error=sign-compare -Wno-deprecated-declarations -Xcompiler=-Wno-deprecated-declarations -Xfatbin=-compress-all -Xcompiler=-fopenmp -MD -MT cuml-cpp/CMakeFiles/cuml++.dir/src/explainer/tree_shap.cu.o -MF cuml-cpp/CMakeFiles/cuml++.dir/src/explainer/tree_shap.cu.o.d -x cu -c /__w/cuml/cuml/cpp/src/explainer/tree_shap.cu -o cuml-cpp/CMakeFiles/cuml++.dir/src/explainer/tree_shap.cu.o
  /__w/cuml/cuml/python/build/cp39-cp39-manylinux_2_28_aarch64/_deps/gputreeshap-src/GPUTreeShap/gpu_treeshap.h(464):
  error #20054-D: dynamic initialization is not supported for a function-scope static __shared__ variable within a __device__/__global__ function
                PathElement<SplitConditionT> s_elements[kBlockSize];
                                             ^
            detected during:
              instantiation of "void gpu_treeshap::detail::ShapKernel<DatasetT,kBlockSize,kRowsPerWarp,SplitConditionT>(DatasetT, size_t, const gpu_treeshap::PathElement<SplitConditionT> *, const size_t *, size_t, double *) [with DatasetT=<unnamed>::DenseDatasetWrapper<float>, kBlockSize=256UL, kRowsPerWarp=1024UL, SplitConditionT=<unnamed>::SplitCondition<float>]" at line 508
              instantiation of "void gpu_treeshap::detail::ComputeShap(DatasetT, const thrust::device_vector<size_t, SizeTAllocatorT> &, const thrust::device_vector<gpu_treeshap::PathElement<SplitConditionT>, PathAllocatorT> &, size_t, double *) [with DatasetT=<unnamed>::DenseDatasetWrapper<float>, SizeTAllocatorT=thrust::device_allocator<size_t>, PathAllocatorT=thrust::device_allocator<gpu_treeshap::PathElement<<unnamed>::SplitCondition<float>>>, SplitConditionT=<unnamed>::SplitCondition<float>]" at line 1313
              instantiation of "void gpu_treeshap::GPUTreeShap(DatasetT, PathIteratorT, PathIteratorT, size_t, PhiIteratorT, PhiIteratorT) [with DeviceAllocatorT=thrust::device_allocator<int>, DatasetT=<unnamed>::DenseDatasetWrapper<float>, PathIteratorT=thrust::detail::normal_iterator<thrust::device_ptr<gpu_treeshap::PathElement<<unnamed>::SplitCondition<float>>>>, PhiIteratorT=thrust::device_ptr<float>]" at line 453 of /__w/cuml/cuml/cpp/src/explainer/tree_shap.cu
              instantiation of "void <unnamed>::gpu_treeshap_impl(ML::Explainer::TreePathInfo<ThresholdT> *, const DataT *, std::size_t, std::size_t, DataT *, std::size_t) [with ThresholdT=float, DataT=float]" at line 791 of /__w/cuml/cuml/cpp/src/explainer/tree_shap.cu

  Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

  1 error detected in the compilation of "/__w/cuml/cuml/cpp/src/explainer/tree_shap.cu".

(build link)

That looks similar to something raft faced with CUDA 12.2: rapidsai/raft#1870.

[jakirkham]

Looks like this needs PR ( rapidsai/gputreeshap#42 )

[jakirkham]

Looks like this needs PR ( rapidsai/gputreeshap#42 )

Testing this in PR ( #5719 )

[jakirkham]

It looks like CUDA 12.2 passes! 🎉

There is a CUDA 11.8 failure with Dask though. After a cursory look, am not seeing any obvious cause

test_sparse_from_dense[8-float32-regularization0-True]

...

E           assert <array_equal: [ 0.64805114  0.40563476 -0.03294854 ... -0.20645222 -0.2884561 -0.13163252] [ 0.64738656  0.40191636 -0.03218034 ... -0.20163251 -0.28334481 -0.12986547] unit_tol=0.005 total_tol=0.0001 with_sign=True>
E            +  where <array_equal: [ 0.64805114  0.40563476 -0.03294854 ... -0.20645222 -0.2884561 -0.13163252] [ 0.64738656  0.40191636 -0.03218034 ... -0.20163251 -0.28334481 -0.12986547] unit_tol=0.005 total_tol=0.0001 with_sign=True> = array_equal(array([ 0.64805114,  0.40563476, -0.03294854,  0.05359353, -0.44779027,\n       -0.20645222, -0.2884561 , -0.13163252], dtype=float32), array([ 0.64738656,  0.40191636, -0.03218034,  0.05311682, -0.45539661,\n       -0.20163251, -0.28334481, -0.12986547]), 0.005, with_sign=True)

test_dask_logistic_regression.py:393: AssertionError

[jameslamb]

there is an 11.8 failure with Dask though

Looks like @dantegd re-triggered that build and now the test passed: https://github.com/rapidsai/cuml/actions/runs/7642195211/job/20839681496?pr=5711

So maybe it is a flaky test? Either way, seems that it should be unrelated to CUDA 12.2 support since it's happening on the 11.8 job.

[vyasr]

Restarting CI now that the pandas 2 fixes are merged in #5758. We can check out the logs tomorrow.

[bdice]

Ready for ops review - CI looks good. Triggering /merge.

[bdice]

/merge