skip old cccl.

Author: trivialfis

python-package/xgboost/testing/multi_target.py

@@ -14,7 +14,7 @@
 
 from .._typing import ArrayLike
 from ..compat import import_cupy
-from ..core import Booster, DMatrix, ExtMemQuantileDMatrix, QuantileDMatrix
+from ..core import Booster, DMatrix, ExtMemQuantileDMatrix, QuantileDMatrix, build_info
 from ..objective import Objective, TreeObjective
 from ..sklearn import XGBClassifier
 from ..training import train
@@ -218,6 +218,10 @@ def run_with_iter(device: Device) -> None:  # pylint: disable=too-many-locals
     X, _, _ = it.as_arrays()
     assert_allclose(device, booster_0.inplace_predict(X), booster_1.inplace_predict(X))
 
+    v = build_info()["THRUST_VERSION"]
+    if v[0] < 3:
+        pytest.xfail("CCCL version too old.")
+
     it = IteratorForTest(
         Xs,
         ys,

src/tree/gpu_hist/leaf_sum.cu

@@ -3,6 +3,7 @@
  */
 #include <thrust/iterator/tabulate_output_iterator.h>  // for make_tabulate_output_iterator
 #include <thrust/scan.h>                               // for inclusive_scan
+#include <thrust/version.h>                            // for THRUST_MAJOR_VERSION
 
 #include <cstddef>                                 // for size_t
 #include <cstdint>                                 // for int32_t
@@ -18,6 +19,12 @@
 #include "xgboost/linalg.h"     // for MatrixView
 #include "xgboost/span.h"       // for Span
 
+#if THRUST_MAJOR_VERSION >= 3
+// do nothing
+#else
+#include "../../common/linalg_op.cuh"  // for tbegin
+#endif
+
 namespace xgboost::tree::cuda_impl {
 void LeafGradSum(Context const* ctx, std::vector<LeafInfo> const& h_leaves,
                  common::Span<GradientQuantiser const> roundings,
@@ -55,8 +62,12 @@ void LeafGradSum(Context const* ctx, std::vector<LeafInfo> const& h_leaves,
       return roundings[t].ToFixedPoint(g);
     });
     // Use an output iterator to implement running sum.
+#if THRUST_MAJOR_VERSION >= 3
     auto out_it = thrust::make_tabulate_output_iterator(
         [=] XGBOOST_DEVICE(std::int32_t idx, GradientPairInt64 v) mutable { out_t(idx) += v; });
+#else
+    auto out_it = linalg::tbegin(out_t);
+#endif
 
     std::size_t n_bytes = 0;
     dh::safe_cuda(cub::DeviceSegmentedReduce::Sum(nullptr, n_bytes, it, out_it, h_leaves.size(),

src/tree/updater_gpu_hist.cuh

@@ -200,7 +200,12 @@ class MultiTargetHistMaker {
     auto d_roundings = this->value_quantizer_->Quantizers();
     // Node indices for all leaves
     std::vector<bst_node_t> leaves_idx(n_leaves);
-
+#if THRUST_MAJOR_VERSION >= 3
+    // do nothing
+#else
+    CHECK_EQ(this->partitioners_.Size(), 1)
+        << "External memory not implemented for old CCCL versions. (thrust < 3.0)";
+#endif
     std::int32_t batch_idx = 0;
     for (auto const& p_part : this->partitioners_) {
       auto leaves = p_part->GetLeaves();