From 5e80c1d2159e00a204ab5db0f5ca3f9ec43187c7 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Fri, 8 Dec 2023 12:03:03 -0600
Subject: [PATCH] Use cuda::proclaim_return_type on device lambda. (#2048)

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: https://github.com/rapidsai/raft/pull/2048
---
 .../raft/spectral/detail/matrix_wrappers.hpp  | 21 +++++++++++--------
 1 file changed, 12 insertions(+), 9 deletions(-)

diff --git a/cpp/include/raft/spectral/detail/matrix_wrappers.hpp b/cpp/include/raft/spectral/detail/matrix_wrappers.hpp
index b90b24e406..6de0dced8a 100644
--- a/cpp/include/raft/spectral/detail/matrix_wrappers.hpp
+++ b/cpp/include/raft/spectral/detail/matrix_wrappers.hpp
@@ -30,6 +30,8 @@
 #include <thrust/reduce.h>
 #include <thrust/system/cuda/execution_policy.h>
 
+#include <cuda/functional>
+
 #include <algorithm>
 
 // =========================================================
@@ -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)