diff --git a/src/common/quantile.cu b/src/common/quantile.cu index 4b110f5e0164..b191de97f6f1 100644 --- a/src/common/quantile.cu +++ b/src/common/quantile.cu @@ -173,13 +173,13 @@ common::Span> MergePath( auto scan_key_it = dh::MakeTransformIterator( thrust::make_counting_iterator(0ul), - [=] __device__(size_t idx) { return dh::SegmentId(out_ptr, idx); }); + [=] XGBOOST_DEVICE(size_t idx) { return dh::SegmentId(out_ptr, idx); }); auto scan_val_it = dh::MakeTransformIterator( - merge_path.data(), [=] __device__(Tuple const &t) -> Tuple { + merge_path.data(), [=] XGBOOST_DEVICE(Tuple const &t) -> Tuple { auto ind = get_ind(t); // == 0 if element is from x // x_counter, y_counter - return thrust::make_tuple(!ind, ind); + return thrust::tuple{!ind, ind}; }); // Compute the index for both x and y (which of the element in a and b are used in each diff --git a/src/data/ellpack_page.cu b/src/data/ellpack_page.cu index c60fe83893bb..d9ea85919bd8 100644 --- a/src/data/ellpack_page.cu +++ b/src/data/ellpack_page.cu @@ -171,11 +171,10 @@ struct WriteCompressedEllpackFunctor { using Tuple = thrust::tuple; __device__ size_t operator()(Tuple out) { - auto e = batch.GetElement(out.get<2>()); + auto e = batch.GetElement(thrust::get<2>(out)); if (is_valid(e)) { // -1 because the scan is inclusive - size_t output_position = - accessor.row_stride * e.row_idx + out.get<1>() - 1; + size_t output_position = accessor.row_stride * e.row_idx + thrust::get<1>(out) - 1; uint32_t bin_idx = 0; if (common::IsCat(feature_types, e.column_idx)) { bin_idx = accessor.SearchBin(e.value, e.column_idx); @@ -192,8 +191,8 @@ template struct TupleScanOp { __device__ Tuple operator()(Tuple a, Tuple b) { // Key equal - if (a.template get<0>() == b.template get<0>()) { - b.template get<1>() += a.template get<1>(); + if (thrust::get<0>(a) == thrust::get<0>(b)) { + thrust::get<1>(b) += thrust::get<1>(a); return b; } // Not equal