ROCM-1254 - fix that functors defined in the cooperative_groups namespace were not usable with user defined types, because they will try to call reduce_sync operations which are only defined with standard types. Add Unit_Thread_Block_Tile_Reduce_Standard_Op_Custom_Type test

Author: g-h-c

projects/clr/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups_reduce.h

@@ -22,6 +22,60 @@ namespace impl {
   using is_param_type_same = __hip_internal::is_same<typename __hip_internal::remove_cvref<T>,
                                                      typename __hip_internal::remove_cvref<U>>;
 
+  template <typename T, typename = void>
+  struct has_add : std::false_type {
+  };
+
+  template <typename T>
+  struct has_add<T,
+                 std::void_t<decltype(__reduce_add_sync<unsigned long long>(0ull, T {}))>
+    > : std::true_type {};
+
+  template <typename T, typename = void>
+  struct has_min : std::false_type {
+  };
+
+  template <typename T>
+  struct has_min<T,
+                 std::void_t<decltype(__reduce_min_sync<unsigned long long>(0ull, T {}))>
+    > : std::true_type {};
+
+  template <typename T, typename = void>
+  struct has_max : std::false_type {
+  };
+
+  template <typename T>
+  struct has_max<T,
+                 std::void_t<decltype(__reduce_max_sync<unsigned long long>(0ull, T {}))>
+    > : std::true_type {};
+
+  template <typename T, typename = void>
+  struct has_and : std::false_type {
+  };
+
+  template <typename T>
+  struct has_and<T,
+                 std::void_t<decltype(__reduce_and_sync<unsigned long long>(0ull, T {}))>
+    > : std::true_type {};
+
+  template <typename T, typename = void>
+  struct has_or : std::false_type {
+  };
+
+  template <typename T>
+  struct has_or<T,
+                 std::void_t<decltype(__reduce_or_sync<unsigned long long>(0ull, T {}))>
+    > : std::true_type {};
+
+  template <typename T, typename = void>
+  struct has_xor : std::false_type {
+  };
+
+  template <typename T>
+  struct has_xor<T,
+                 std::void_t<decltype(__reduce_xor_sync<unsigned long long>(0ull, T {}))>
+    > : std::true_type {};
+
   // we can call reduce() only the block tiles that have a compile-time size
   template <class TyGroup>
   struct isTiledGroup : __hip_internal::false_type {
@@ -69,17 +123,23 @@ __CG_QUALIFIER__ auto reduce(const TyGroup& group, TyVal&& val, TyFn&& op) -> de
   // need to apply the active mask
   mask &= __activemask();
 
-  if constexpr (__hip_internal::is_same<Op, cooperative_groups::plus<Val>>::value) {
+  if constexpr (__hip_internal::is_same<Op, cooperative_groups::plus<Val>>::value &&
+                impl::has_add<Val>::value) {
     return __reduce_add_sync(mask, val);
-  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::less<Val>>::value) {
+  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::less<Val>>::value &&
+                impl::has_min<Val>::value) {
     return __reduce_min_sync(mask, val);
-  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::greater<Val>>::value) {
+  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::greater<Val>>::value &&
+                impl::has_max<Val>::value) {
     return __reduce_max_sync(mask, val);
-  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::bit_and<Val>>::value) {
+  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::bit_and<Val>>::value &&
+                impl::has_and<Val>::value) {
     return __reduce_and_sync(mask, val);
-  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::bit_or<Val>>::value) {
+  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::bit_or<Val>>::value &&
+                impl::has_or<Val>::value) {
     return __reduce_or_sync(mask, val);
-  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::bit_xor<Val>>::value) {
+  } else if constexpr (__hip_internal::is_same<Op, cooperative_groups::bit_xor<Val>>::value &&
+                impl::has_xor<Val>::value) {
     return __reduce_xor_sync(mask, val);
   } else {
     return __reduce_op_sync(mask, val, op, nullptr);

projects/hip-tests/catch/config/configs/unit/cooperativeGrps.yaml

@@ -242,6 +242,8 @@ cooperativeGrps:
     <<: *level_2
     # Rock_Window_Failures_on_gfx1151
     disabled: [amd_windows]
+  Unit_Thread_Block_Tile_Reduce_Standard_Op_Custom_Type:
+    <<: *level_2
   Unit_Thread_Block_Tile_Reduce_Trivially_Copyable_Parameters:
     <<: *level_2
     # Rock_Window_Failures_on_gfx1151

projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc

@@ -1000,6 +1000,54 @@ TEST_CASE(Unit_Thread_Block_Tile_Reduce_All_Parameter_Sizes)
   }
 }
 
+struct Point {
+    int x;
+    int y;
+
+    __device__ Point operator+(const Point& rhs)
+    {
+      return { x + rhs.x, y + rhs.y };
+    }
+
+};
+
+__global__ void sumPoints(Point* result)
+{
+  cg::thread_block mygroup = cg::this_thread_block();
+  auto mytile = cg::tiled_partition<32>(mygroup);
+  Point input;
+  
+  input.x = threadIdx.x;
+  input.y = threadIdx.x;
+
+  __syncwarp();
+  *result = cg::reduce(mytile, input, cooperative_groups::plus<Point> {});
+}
+
+// using a standard functor in the cooperative_groups namespace with a type that is not primitive
+TEST_CASE(Unit_Thread_Block_Tile_Reduce_Standard_Op_Custom_Type)
+{
+  LinearAllocGuard<Point> h_result(LinearAllocs::malloc, sizeof(Point) * 32);
+  LinearAllocGuard<Point> d_result(LinearAllocs::hipMalloc, sizeof(Point) * 32);
+  dim3 gridDim = { 1 };
+  dim3 blockDim = { 32 };
+  void* devicePtr = d_result.ptr();
+  void* args[] = { &devicePtr };
+  int expected = 31 * 16;
+
+  HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(sumPoints), gridDim, blockDim, args, 0, nullptr));
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipGetLastError());
+  HIP_CHECK(hipMemcpy(h_result.host_ptr(), d_result.ptr(),
+                      h_result.size_bytes(), hipMemcpyDeviceToHost));
+
+  for (int i = 0; i < 32; i++) {
+    INFO("Expected x: " << expected << " got: " << *h_result.host_ptr());
+    INFO("Expected y: " << expected << " got: " << *h_result.host_ptr());
+    REQUIRE((h_result.host_ptr()->x == expected && h_result.host_ptr()->y == expected));
+  }
+}
+
 TEMPLATE_TEST_CASE(Unit_Thread_Block_Coalesced_Reduce_arithmetic, int, unsigned int, long long,
                    unsigned long long, float, half, double)
 {