From 053fda1ec72cecbdf75426babc6ef39fef4b81a3 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Sat, 17 Sep 2022 12:24:01 +0200 Subject: [PATCH] fix CUDA 9.2 (this time for real) --- common/cuda_hip/matrix/csr_kernels.hpp.inc | 14 ++++---------- 1 file changed, 4 insertions(+), 10 deletions(-) diff --git a/common/cuda_hip/matrix/csr_kernels.hpp.inc b/common/cuda_hip/matrix/csr_kernels.hpp.inc index 223574261f8..85b7798cb5d 100644 --- a/common/cuda_hip/matrix/csr_kernels.hpp.inc +++ b/common/cuda_hip/matrix/csr_kernels.hpp.inc @@ -1122,17 +1122,11 @@ void fallback_transpose(std::shared_ptr exec, out_col_idxs); exec->copy(nnz, in_vals, out_vals); exec->copy(nnz, in_col_idxs, out_row_idxs.get_data()); - auto zip_it = thrust::make_zip_iterator( - thrust::make_tuple(thrust::device_pointer_cast(out_row_idxs.get_data()), - thrust::device_pointer_cast(out_col_idxs), - thrust::device_pointer_cast(out_vals))); + auto loc_it = thrust::make_zip_iterator( + thrust::make_tuple(out_row_idxs.get_data(), out_col_idxs)); using tuple_type = - thrust::tuple>; - thrust::sort(thrust::device, zip_it, zip_it + nnz, - [] __device__(const tuple_type& a, const tuple_type& b) { - return thrust::tie(thrust::get<0>(a), thrust::get<1>(a)) < - thrust::tie(thrust::get<0>(b), thrust::get<1>(b)); - }); + thrust::tuple>; + thrust::sort_by_key(thrust::device, loc_it, loc_it + nnz, out_vals); components::convert_idxs_to_ptrs(exec, out_row_idxs.get_data(), nnz, out_num_rows, out_row_ptrs); }