Fix non default constructible input types test for cub::FindIf (#7447)

* Fix non default constructible input types test for cub::FindIf

- make NotDefaultConstructible non default constructible by removing init in input param
- make iterator raw so that vectorized path is taken

* Use thrust::tabulate to initialized input array and other nits

* Make functor non generic

* Shortcircuit nigation train

Author: gonidelis

cub/test/catch2_test_device_find_if.cu

@@ -6,6 +6,7 @@
 #include <cub/device/device_find.cuh>
 
 #include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/tabulate.h>
 
 #include <cuda/iterator>
 
@@ -207,7 +208,7 @@ struct NotDefaultConstructible
 {
   int value_;
 
-  __host__ __device__ constexpr explicit NotDefaultConstructible(int value = 0)
+  __host__ __device__ constexpr explicit NotDefaultConstructible(int value)
       : value_(value)
   {}
 
@@ -228,18 +229,21 @@ struct NotDefaultConstructible
   }
 };
 
-struct converter
+struct index_to_value
 {
-  __host__ __device__ constexpr NotDefaultConstructible operator()(const int val) const noexcept
+  __host__ __device__ NotDefaultConstructible operator()(int i)
   {
-    return NotDefaultConstructible{val};
+    return NotDefaultConstructible{static_cast<int>(i)};
   }
 };
 
+static_assert(!cuda::std::is_default_constructible_v<NotDefaultConstructible>,
+              "NotDefaultConstructible should not be default constructible");
+
 C2H_TEST("Device find_if works with non default constructible types", "[device][find_if]")
 {
-  using input_t  = int32_t;
-  using offset_t = int32_t;
+  using input_t  = NotDefaultConstructible;
+  using offset_t = int;
 
   constexpr offset_t min_items = 1;
   constexpr offset_t max_items = 10'000; // 10k items for reasonable test time
@@ -251,16 +255,19 @@ C2H_TEST("Device find_if works with non default constructible types", "[device][
       min_items,
       max_items,
     }));
-  const input_t val_to_find = static_cast<input_t>(num_items - 1);
+  const auto val_to_find = static_cast<int>(num_items - 1);
 
   CAPTURE(num_items, val_to_find);
 
-  // counting_iterator input
-  auto c_it = cuda::make_transform_iterator(cuda::make_counting_iterator(input_t{0}), converter{});
-  {
-    c2h::device_vector<offset_t> out_result(1, thrust::no_init);
-    auto predicate = thrust::detail::equal_to_value<NotDefaultConstructible>{NotDefaultConstructible{val_to_find}};
-    find_if(c_it, thrust::raw_pointer_cast(out_result.data()), predicate, num_items);
-    REQUIRE(val_to_find == out_result[0]);
-  }
+  // raw device iterator to some device vector so that vectorized path is taken
+  c2h::device_vector<input_t> d_vec(num_items, NotDefaultConstructible(0));
+
+  // fill with arbitrary values dont use c2h gen because NotDefaultConstructible is not default constructible
+  thrust::tabulate(c2h::device_policy, d_vec.begin(), d_vec.end(), index_to_value{});
+
+  auto it = thrust::raw_pointer_cast(d_vec.data());
+  c2h::device_vector<offset_t> out_result(1, thrust::no_init);
+  auto predicate = thrust::detail::equal_to_value<NotDefaultConstructible>{NotDefaultConstructible{val_to_find}};
+  find_if(it, thrust::raw_pointer_cast(out_result.data()), predicate, num_items);
+  REQUIRE(val_to_find == out_result[0]);
 }