Skip to content

Commit

Permalink
Use cuda::proclaim_return_type on device lambda. (#2048)
Browse files Browse the repository at this point in the history
This PR is needed to change the one piece of RAFT that requires `cuda::proclaim_return_type` for compatibility with CCCL (Thrust) 2.2.0. This pulls out part of the diff of #1464, which we will be able to close in favor of a new PR after this is merged.

Authors:
  - Bradley Dice (https://github.com/bdice)

Approvers:
  - Divye Gala (https://github.com/divyegala)

URL: #2048
  • Loading branch information
bdice authored Dec 8, 2023
1 parent e999100 commit 5e80c1d
Showing 1 changed file with 12 additions and 9 deletions.
21 changes: 12 additions & 9 deletions cpp/include/raft/spectral/detail/matrix_wrappers.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,8 @@
#include <thrust/reduce.h>
#include <thrust/system/cuda/execution_policy.h>

#include <cuda/functional>

#include <algorithm>

// =========================================================
Expand Down Expand Up @@ -107,15 +109,16 @@ class vector_t {

value_type nrm1() const
{
return thrust::reduce(thrust_policy,
buffer_.data(),
buffer_.data() + buffer_.size(),
value_type{0},
[] __device__(auto left, auto right) {
auto abs_left = left > 0 ? left : -left;
auto abs_right = right > 0 ? right : -right;
return abs_left + abs_right;
});
return thrust::reduce(
thrust_policy,
buffer_.data(),
buffer_.data() + buffer_.size(),
value_type{0},
cuda::proclaim_return_type<value_type>([] __device__(auto left, auto right) {
auto abs_left = left > 0 ? left : -left;
auto abs_right = right > 0 ? right : -right;
return abs_left + abs_right;
}));
}

void fill(value_type value)
Expand Down

0 comments on commit 5e80c1d

Please sign in to comment.