[REVIEW] Add a double type dispatcher.

CHANGELOG.md

@@ -48,6 +48,7 @@
 - PR #5658 Add `filter_tokens` nvtext API
 - PR #5666 Add `filter_characters_of_type` strings API
 - PR #5673 Always build and test with per-thread default stream enabled in the GPU CI build
+- PR #5716 Add `double_type_dispatcher` to libcudf
 - PR #5572 Add `cudf::encode` API.
 
 ## Improvements

cpp/include/cudf/utilities/type_dispatcher.hpp

@@ -419,5 +419,58 @@ CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) type_dispatcher(cudf::data_ty
   }
 }
 
+namespace detail {
+template <typename T1>
+struct double_type_dispatcher_second_type {
+#pragma nv_exec_check_disable
+  template <typename T2, typename F, typename... Ts>
+  CUDA_HOST_DEVICE_CALLABLE decltype(auto) operator()(F&& f, Ts&&... args) const
+  {
+    return f.template operator()<T1, T2>(std::forward<Ts>(args)...);
+  }
+};
+
+template <template <cudf::type_id> typename IdTypeMap>
+struct double_type_dispatcher_first_type {
+#pragma nv_exec_check_disable
+  template <typename T1, typename F, typename... Ts>
+  CUDA_HOST_DEVICE_CALLABLE decltype(auto) operator()(cudf::data_type type2,
+                                                      F&& f,
+                                                      Ts&&... args) const
+  {
+    return type_dispatcher<IdTypeMap>(type2,
+                                      detail::double_type_dispatcher_second_type<T1>{},
+                                      std::forward<F>(f),
+                                      std::forward<Ts>(args)...);
+  }
+};
+}  // namespace detail
+
+/**
+ * @brief Dispatches two type template parameters to a callable.
+ *
+ * This function expects a callable `f` with an `operator()` template accepting
+ * two typename template parameters `T1` and `T2`.
+ *
+ * @param type1 The `data_type` used to dispatch a type for the first template
+ * parameter of the callable `F`
+ * @param type2 The `data_type` used to dispatch a type for the second template
+ * parameter of the callable `F`
+ * @param args Parameter pack forwarded to the `operator()` invocation `F`.
+ */
+#pragma nv_exec_check_disable
+template <template <cudf::type_id> typename IdTypeMap = id_to_type_impl, typename F, typename... Ts>
+CUDA_HOST_DEVICE_CALLABLE constexpr decltype(auto) double_type_dispatcher(cudf::data_type type1,
+                                                                          cudf::data_type type2,
+                                                                          F&& f,
+                                                                          Ts&&... args)
+{
+  return type_dispatcher<IdTypeMap>(type1,
+                                    detail::double_type_dispatcher_first_type<IdTypeMap>{},
+                                    type2,
+                                    std::forward<F>(f),
+                                    std::forward<Ts>(args)...);
+}
+
 /** @} */  // end of group
 }  // namespace cudf

cpp/tests/types/type_dispatcher_test.cu

@@ -84,4 +84,95 @@ TEST_P(IdDispatcherTest, IdToType)
   EXPECT_TRUE(cudf::type_dispatcher(cudf::data_type{t}, verify_dispatched_type{}, t));
 }
 
+template <typename T>
+struct DoubleTypedDispatcherTest : public DispatcherTest {
+};
+
+TYPED_TEST_CASE(DoubleTypedDispatcherTest, cudf::test::AllTypes);
+
+namespace {
+template <typename Expected1, typename Expected2>
+struct two_type_tester {
+  template <typename Dispatched1, typename Dispatched2>
+  bool operator()()
+  {
+    return std::is_same<Expected1, Dispatched1>::value &&
+           std::is_same<Expected2, Dispatched2>::value;
+  }
+};
+}  // namespace
+
+TYPED_TEST(DoubleTypedDispatcherTest, TypeToId)
+{
+  EXPECT_TRUE(cudf::double_type_dispatcher(cudf::data_type{cudf::type_to_id<TypeParam>()},
+                                           cudf::data_type{cudf::type_to_id<TypeParam>()},
+                                           two_type_tester<TypeParam, TypeParam>{}));
+}
+
+namespace {
+struct verify_double_dispatched_type {
+  template <typename T1, typename T2>
+  __host__ __device__ bool operator()(cudf::type_id id1, cudf::type_id id2)
+  {
+    return id1 == cudf::type_to_id<T1>() && id2 == cudf::type_to_id<T2>();
+  }
+};
+
+__global__ void double_dispatch_test_kernel(cudf::type_id id1, cudf::type_id id2, bool* d_result)
+{
+  if (0 == threadIdx.x + blockIdx.x * blockDim.x)
+    *d_result = cudf::double_type_dispatcher(
+      cudf::data_type{id1}, cudf::data_type{id2}, verify_double_dispatched_type{}, id1, id2);
+}
+}  // namespace
+
+TYPED_TEST(DoubleTypedDispatcherTest, DeviceDoubleDispatch)
+{
+  thrust::device_vector<bool> result(1, false);
+  double_dispatch_test_kernel<<<1, 1>>>(
+    cudf::type_to_id<TypeParam>(), cudf::type_to_id<TypeParam>(), result.data().get());
+  CUDA_TRY(cudaDeviceSynchronize());
+  EXPECT_EQ(true, result[0]);
+}
+
+struct IdDoubleDispatcherTest : public DispatcherTest,
+                                public testing::WithParamInterface<cudf::type_id> {
+};
+
+INSTANTIATE_TEST_CASE_P(TestAllIds,
+                        IdDoubleDispatcherTest,
+                        testing::ValuesIn(cudf::test::all_type_ids));
+
+TEST_P(IdDoubleDispatcherTest, IdToType)
+{
+  // Test double-dispatch of all types using the same type for both dispatches
+  auto t = GetParam();
+  EXPECT_TRUE(cudf::double_type_dispatcher(
+    cudf::data_type{t}, cudf::data_type{t}, verify_double_dispatched_type{}, t, t));
+}
+
+struct IdFixedDoubleDispatcherTest : public DispatcherTest,
+                                     public testing::WithParamInterface<cudf::type_id> {
+};
+
+INSTANTIATE_TEST_CASE_P(TestAllIds,
+                        IdFixedDoubleDispatcherTest,
+                        testing::ValuesIn(cudf::test::all_type_ids));
+
+TEST_P(IdFixedDoubleDispatcherTest, IdToType)
+{
+  // Test double-dispatch of all types against one fixed type, in each direction
+  auto t = GetParam();
+  EXPECT_TRUE(cudf::double_type_dispatcher(cudf::data_type{t},
+                                           cudf::data_type{cudf::type_to_id<float>()},
+                                           verify_double_dispatched_type{},
+                                           t,
+                                           cudf::type_to_id<float>()));
+  EXPECT_TRUE(cudf::double_type_dispatcher(cudf::data_type{cudf::type_to_id<float>()},
+                                           cudf::data_type{t},
+                                           verify_double_dispatched_type{},
+                                           cudf::type_to_id<float>(),
+                                           t));
+}
+
 CUDF_TEST_PROGRAM_MAIN()