// SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. All rights reserved.
// SPDX-License-Identifier: BSD-3-Clause

#include <cub/util_macro.cuh>
#include <cub/warp/warp_reduce.cuh>

#include <cuda/cmath>
#include <cuda/functional>
#include <cuda/iterator>
#include <cuda/ptx>
#include <cuda/std/__floating_point/arithmetic.h>
#include <cuda/std/__floating_point/cast.h>
#include <cuda/std/__functional/invoke.h>
#include <cuda/std/functional>
#include <cuda/std/limits>
#include <cuda/std/type_traits>

#include <array>
#include <numeric>

#include <test_util.h>

#include "cub_test_macros.h"
#include <c2h/check_results.cuh>
#include <c2h/custom_type.h>
#include <c2h/operator.cuh>

/***********************************************************************************************************************
 * Constants
 **********************************************************************************************************************/

inline constexpr int warp_size            = 32;
inline constexpr auto total_warps         = 4u;
inline constexpr int num_items_per_thread = 4;

/***********************************************************************************************************************
 * Kernel
 **********************************************************************************************************************/

template <unsigned LogicalWarpThreads, bool EnableNumItems = false, typename T, typename Output, typename ReductionOp>
__device__ void warp_reduce_function(T& thread_data, Output* output, ReductionOp reduction_op, int num_items = 0)
{
  using warp_reduce_t = cub::WarpReduce<Output, LogicalWarpThreads>;
  using storage_t     = typename warp_reduce_t::TempStorage;
  __shared__ storage_t storage[total_warps];
  constexpr bool is_power_of_two = cuda::is_power_of_two(LogicalWarpThreads);
  auto lane                      = cuda::ptx::get_sreg_laneid();
  auto logical_warp              = is_power_of_two ? threadIdx.x / LogicalWarpThreads : threadIdx.x / warp_size;
  auto logical_lane              = is_power_of_two ? threadIdx.x % LogicalWarpThreads : lane;
  auto limit                     = EnableNumItems ? num_items : LogicalWarpThreads;
  if (!is_power_of_two && lane >= limit)
  {
    return;
  }
  const warp_reduce_t warp_reduce{storage[logical_warp]};
  using result_t = decltype(reduction_op(warp_reduce, thread_data));
  result_t result;
  if constexpr (EnableNumItems)
  {
    result = reduction_op(warp_reduce, thread_data, num_items);
  }
  else
  {
    result = reduction_op(warp_reduce, thread_data);
  }
  if (logical_lane == 0)
  {
    output[logical_warp] = result;
  }
}

template <unsigned LogicalWarpThreads, bool EnableNumItems, typename T, typename ReductionOp>
__global__ void warp_reduce_kernel(T* input, T* output, ReductionOp reduction_op, int num_items = 0)
{
  auto thread_data = input[threadIdx.x];
  warp_reduce_function<LogicalWarpThreads, EnableNumItems>(thread_data, output, reduction_op, num_items);
}

template <unsigned LogicalWarpThreads, typename T, typename ReductionOp>
__global__ void warp_reduce_multiple_items_kernel(T* input, T* output, ReductionOp reduction_op)
{
  T thread_data[num_items_per_thread];
  for (int i = 0; i < num_items_per_thread; ++i)
  {
    thread_data[i] = input[threadIdx.x * num_items_per_thread + i];
  }
  warp_reduce_function<LogicalWarpThreads>(thread_data, output, reduction_op);
}

template <typename Op, typename T>
struct warp_reduce_t
{
  template <int LogicalWarpThreads>
  __device__ auto operator()(cub::WarpReduce<T, LogicalWarpThreads> warp_reduce, T& data) const
  {
    return warp_reduce.Reduce(data, Op{});
  }

  template <int LogicalWarpThreads>
  __device__ auto operator()(cub::WarpReduce<T, LogicalWarpThreads> warp_reduce, T& data, int num_items) const
  {
    return warp_reduce.Reduce(data, Op{}, num_items);
  }
};

template <typename T>
struct warp_reduce_t<cuda::std::plus<>, T>
{
  template <int LogicalWarpThreads, typename... TArgs>
  __device__ auto operator()(cub::WarpReduce<T, LogicalWarpThreads> warp_reduce, TArgs&&... args) const
  {
    return warp_reduce.Sum(args...);
  }
};

template <typename T>
struct warp_reduce_t<cuda::maximum<>, T>
{
  template <int LogicalWarpThreads, typename... TArgs>
  __device__ auto operator()(cub::WarpReduce<T, LogicalWarpThreads> warp_reduce, TArgs&&... args) const
  {
    return warp_reduce.Max(args...);
  }
};

template <typename T>
struct warp_reduce_t<cuda::minimum<>, T>
{
  template <int LogicalWarpThreads, typename... TArgs>
  __device__ auto operator()(cub::WarpReduce<T, LogicalWarpThreads> warp_reduce, TArgs&&... args) const
  {
    return warp_reduce.Min(args...);
  }
};

template <int LogicalWarpThreads, bool EnableNumItems = false, typename T, typename... TArgs>
void warp_reduce_launch(c2h::device_vector<T>& input, c2h::device_vector<T>& output, TArgs... args)
{
  warp_reduce_kernel<LogicalWarpThreads, EnableNumItems><<<1, total_warps * warp_size>>>(
    thrust::raw_pointer_cast(input.data()), thrust::raw_pointer_cast(output.data()), args...);

  REQUIRE(cudaSuccess == cudaPeekAtLastError());
  REQUIRE(cudaSuccess == cudaDeviceSynchronize());
}

template <int LogicalWarpThreads, typename T, typename... TArgs>
void warp_reduce_multiple_items_launch(c2h::device_vector<T>& input, c2h::device_vector<T>& output, TArgs... args)
{
  warp_reduce_multiple_items_kernel<LogicalWarpThreads><<<1, total_warps * warp_size>>>(
    thrust::raw_pointer_cast(input.data()), thrust::raw_pointer_cast(output.data()), args...);

  REQUIRE(cudaSuccess == cudaPeekAtLastError());
  REQUIRE(cudaSuccess == cudaDeviceSynchronize());
}

/***********************************************************************************************************************
 * Types
 **********************************************************************************************************************/

using custom_t =
  c2h::custom_type_t<c2h::accumulateable_t, c2h::equal_comparable_t, c2h::lexicographical_less_comparable_t>;

using full_type_list =
  c2h::type_list<uint8_t,
                 uint16_t,
                 int32_t,
                 int64_t,
                 custom_t,
#if _CCCL_CTK_AT_LEAST(13, 0)
                 ulonglong4_16a,
#else // _CCCL_CTK_AT_LEAST(13, 0)
                 ulonglong4,
#endif // _CCCL_CTK_AT_LEAST(13, 0)
                 uchar3,
                 short2>;

using builtin_type_list = c2h::type_list<uint8_t, uint16_t, int32_t, int64_t>;

// clang-format off
using floating_point_redux_type_list = c2h::type_list<
float
, double
#if TEST_HALF_T()
, __half
#endif // TEST_HALF_T()
#if TEST_BF_T()
, __nv_bfloat16
#endif // TEST_BF_T()
>;
// clang-format on

using predefined_op_list = c2h::type_list<cuda::std::plus<>, cuda::maximum<>, cuda::minimum<>>;

using predefined_min_max_op_list = c2h::type_list<cuda::maximum<>, cuda::minimum<>>;

// clang-format off
using bitwise_type_list = c2h::type_list<
  int32_t,
  uint32_t,
  int64_t,
  uint64_t
#if TEST_INT128()
  , __int128_t
  , __uint128_t
#endif // TEST_INT128()
>;

using wide_integer_type_list = c2h::type_list<
  int64_t,
  uint64_t
#if TEST_INT128()
  , __int128_t
  , __uint128_t
#endif // TEST_INT128()
>;
// clang-format on

using bitwise_op_list = c2h::type_list<cuda::std::bit_and<>, cuda::std::bit_or<>, cuda::std::bit_xor<>>;

using bool_op_list = c2h::type_list<cuda::std::logical_or<>, cuda::std::logical_and<>>;

using logical_warp_threads = c2h::enum_type_list<unsigned, 32, 16, 9, 7, 1>;

/***********************************************************************************************************************
 * Reference
 **********************************************************************************************************************/

_CCCL_DIAG_PUSH
_CCCL_DIAG_SUPPRESS_MSVC(4244) // numeric(33): C: '=': conversion from 'int' to '_Ty', possible loss of data

template <typename PredefinedOp, typename T>
void compute_host_reference(
  const c2h::host_vector<T>& h_in,
  c2h::host_vector<T>& h_out,
  int logical_warps,
  int logical_warp_threads,
  int items_per_logical_warp = 0,
  int items_per_thread       = 1)
{
  const auto identity    = identity_v<PredefinedOp, T>;
  items_per_logical_warp = items_per_logical_warp == 0 ? logical_warp_threads : items_per_logical_warp;
  for (unsigned i = 0; i < total_warps; ++i)
  {
    for (int j = 0; j < logical_warps; ++j)
    {
      auto start =
        h_in.begin()
        + (i * warp_size + j * logical_warp_threads) * items_per_thread; // NOLINT(bugprone-misplaced-widening-cast)
      auto end = start + static_cast<long>(items_per_logical_warp) * items_per_thread;
      // NOLINTNEXTLINE(bugprone-misplaced-widening-cast)
      h_out[i * logical_warps + j] = static_cast<T>(std::accumulate(start, end, identity, PredefinedOp{}));
    }
  }
}

_CCCL_DIAG_POP

std::array<unsigned, 3> get_test_config(unsigned logical_warp_threads, unsigned items_per_thread = 1)
{
  const bool is_power_of_two = cuda::is_power_of_two(logical_warp_threads);
  auto logical_warps         = is_power_of_two ? warp_size / logical_warp_threads : 1;
  auto input_size            = total_warps * warp_size * items_per_thread;
  auto output_size           = total_warps * logical_warps;
  return {input_size, output_size, logical_warps};
}

/***********************************************************************************************************************
 * Test cases
 **********************************************************************************************************************/

CUB_TEST("WarpReduce::Sum, full_type_list",
         "[reduce][warp][predefined_op][full]",
         CUB_SMALL,
         full_type_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<1, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  CAPTURE(c2h::type_name<T>(), c2h::type_name<T>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<cuda::std::plus<>, T>{});

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<cuda::std::plus<>>(h_in, h_out, logical_warps, logical_warp_threads);
  verify_results(h_out, d_out);
}

CUB_TEST("WarpReduce::Reduce, wide integers",
         "[reduce][warp][predefined_op][redux]",
         CUB_SMALL,
         wide_integer_type_list,
         predefined_op_list)
{
  using T                            = c2h::get<0, TestType>;
  using predefined_op                = c2h::get<1, TestType>;
  constexpr int logical_warp_threads = 32;
  const int valid_items              = GENERATE(1, 17, 31, 32);
  constexpr auto input_size          = total_warps * warp_size;
  constexpr auto output_size         = total_warps;
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), valid_items);

  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads, true>(d_in, d_out, warp_reduce_t<predefined_op, T>{}, valid_items);

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, 1, logical_warp_threads, valid_items);
  verify_results(h_out, d_out);
}

CUB_TEST("WarpReduce::Sum/Max/Min, builtin types",
         "[reduce][warp][predefined_op][full]",
         CUB_SMALL,
         builtin_type_list,
         predefined_op_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  using predefined_op                           = c2h::get<1, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<2, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<predefined_op, T>{});

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads);
  verify_results(h_out, d_out);
}

CUB_TEST("WarpReduce::Max/Min, floating-point redux types",
         "[reduce][warp][predefined_op][redux]",
         CUB_SMALL,
         floating_point_redux_type_list,
         predefined_min_max_op_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  using predefined_op                           = c2h::get<1, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<2, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<predefined_op, T>{});

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads);
  if constexpr (cuda::std::is_same_v<T, float>)
  {
    verify_results_exact(h_out, d_out);
  }
  else
  {
    verify_results(h_out, d_out);
  }
}

CUB_TEST("WarpReduce::Max/Min, floating-point NaN inputs are ignored",
         "[reduce][warp][predefined_op][redux]",
         CUB_SMALL,
         floating_point_redux_type_list,
         predefined_min_max_op_list)
{
  using T               = c2h::get<0, TestType>;
  using predefined_op   = c2h::get<1, TestType>;
  constexpr bool is_min = cuda::std::is_same_v<predefined_op, cuda::minimum<>>;
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>());
  // -NaN and +NaN are the lowest and highest values
  const auto nan        = cuda::std::numeric_limits<T>::quiet_NaN();
  const auto nan_per_op = is_min ? cuda::std::__fp_neg(nan) : nan;
  c2h::host_vector<T> h_in(total_warps * warp_size);
  for (unsigned i = 0; i < h_in.size(); ++i)
  {
    const auto lane = i % warp_size;
    h_in[i]         = lane == 0 ? nan_per_op : cuda::std::__fp_cast<T>(static_cast<float>(lane));
  }
  c2h::device_vector<T> d_in = h_in;
  c2h::device_vector<T> d_out(total_warps);
  warp_reduce_launch<warp_size>(d_in, d_out, warp_reduce_t<predefined_op, T>{});

  const auto expected = is_min ? 1.0f : warp_size - 1.0f;
  const c2h::host_vector<T> h_expected(total_warps, cuda::std::__fp_cast<T>(expected));
  verify_results_exact(h_expected, d_out);
}

#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only
CUB_TEST("WarpReduce::Reduce, bitwise types",
         "[reduce][warp][predefined_op][redux]",
         CUB_SMALL,
         bitwise_type_list,
         bitwise_op_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  using predefined_op                           = c2h::get<1, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<2, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<predefined_op, T>{});

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads);
  verify_results_exact(h_out, d_out);
}

#endif // _CCCL_COMPILER(GCC, >=, 8)

CUB_TEST(
  "WarpReduce::Reduce, bool", "[reduce][warp][predefined_op][vote]", CUB_SMALL, bool_op_list, logical_warp_threads)
{
  using predefined_op                           = c2h::get<0, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<1, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  const unsigned valid_items = GENERATE_COPY(logical_warp_threads, take(2, random(1u, logical_warp_threads)));
  // sparse inputs exercise __any_sync (true is rare), dense inputs exercise __all_sync (false is rare)
  const bool sparse = GENERATE(false, true);
  CAPTURE(c2h::type_name<predefined_op>(), logical_warp_threads, valid_items, sparse);

  c2h::device_vector<uint8_t> d_rand(input_size);
  c2h::gen(C2H_SEED(5), d_rand);
  const c2h::host_vector<uint8_t> h_rand = d_rand;
  c2h::host_vector<bool> h_in(input_size);

  // generate true values with probability 1/logical_warp_threads
  constexpr auto uint8_total    = cuda::std::numeric_limits<uint8_t>::max() + 1;
  constexpr auto true_threshold = cuda::ceil_div(uint8_total, logical_warp_threads);
  for (unsigned i = 0; i < input_size; ++i)
  {
    h_in[i] = (h_rand[i] < true_threshold) == sparse;
  }
  c2h::device_vector<bool> d_in = h_in;
  c2h::device_vector<bool> d_out(output_size);
  if (valid_items == logical_warp_threads)
  {
    warp_reduce_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<predefined_op, bool>{});
  }
  else
  {
    warp_reduce_launch<logical_warp_threads, true>(
      d_in, d_out, warp_reduce_t<predefined_op, bool>{}, static_cast<int>(valid_items));
  }

  c2h::host_vector<bool> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads, valid_items);
  verify_results(h_out, d_out);
}

CUB_TEST("WarpReduce::CustomSum", "[reduce][warp][generic][full]", CUB_SMALL, full_type_list, logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<1, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  CAPTURE(c2h::type_name<T>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(1), d_in);
  warp_reduce_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<custom_plus, T>{});

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<cuda::std::plus<>>(h_in, h_out, logical_warps, logical_warp_threads);
  verify_results(h_out, d_out);
}

//----------------------------------------------------------------------------------------------------------------------
// partial

CUB_TEST("WarpReduce::Sum/Max/Min Partial",
         "[reduce][warp][predefined_op][partial]",
         CUB_SMALL,
         builtin_type_list,
         predefined_op_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  using predefined_op                           = c2h::get<1, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<2, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  const int valid_items                         = GENERATE_COPY(take(2, random(1u, logical_warp_threads)));
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), logical_warp_threads, valid_items);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads, true>(d_in, d_out, warp_reduce_t<predefined_op, T>{}, valid_items);

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads, valid_items);
  verify_results(h_out, d_out);
}

#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only

CUB_TEST("WarpReduce::Reduce, bitwise types, partial",
         "[reduce][warp][predefined_op][redux][partial]",
         CUB_SMALL,
         bitwise_type_list,
         bitwise_op_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  using predefined_op                           = c2h::get<1, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<2, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  const int valid_items                         = GENERATE_COPY(take(2, random(1u, logical_warp_threads)));
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), logical_warp_threads, valid_items);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads, true>(d_in, d_out, warp_reduce_t<predefined_op, T>{}, valid_items);

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads, valid_items);
  verify_results_exact(h_out, d_out);
}

#endif // _CCCL_COMPILER(GCC, >=, 8)

CUB_TEST("WarpReduce::Sum", "[reduce][warp][generic][partial]", CUB_SMALL, full_type_list, logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<1, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads);
  const int valid_items                         = GENERATE_COPY(take(2, random(1u, logical_warp_threads)));
  CAPTURE(c2h::type_name<T>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_launch<logical_warp_threads, true>(d_in, d_out, warp_reduce_t<custom_plus, T>{}, valid_items);

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<cuda::std::plus<>>(h_in, h_out, logical_warps, logical_warp_threads, valid_items);
  verify_results(h_out, d_out);
}

//----------------------------------------------------------------------------------------------------------------------
// multiple items per thread

CUB_TEST("WarpReduce::Sum/Max/Min Multiple Items Per Thread",
         "[reduce][warp][predefined_op][full]",
         CUB_SMALL,
         builtin_type_list,
         predefined_op_list,
         logical_warp_threads)
{
  using T                                       = c2h::get<0, TestType>;
  using predefined_op                           = c2h::get<1, TestType>;
  constexpr auto logical_warp_threads           = c2h::get<2, TestType>::value;
  auto [input_size, output_size, logical_warps] = get_test_config(logical_warp_threads, num_items_per_thread);
  CAPTURE(c2h::type_name<T>(), c2h::type_name<predefined_op>(), logical_warp_threads);
  c2h::device_vector<T> d_in(input_size);
  c2h::device_vector<T> d_out(output_size);
  c2h::gen(C2H_SEED(10), d_in);
  warp_reduce_multiple_items_launch<logical_warp_threads>(d_in, d_out, warp_reduce_t<predefined_op, T>{});

  const c2h::host_vector<T> h_in = d_in;
  c2h::host_vector<T> h_out(output_size);
  compute_host_reference<predefined_op>(h_in, h_out, logical_warps, logical_warp_threads, 0, num_items_per_thread);
  verify_results(h_out, d_out);
}