[SYCLomatic] Refine migration for thrust::max , thrust::min and add them to api-query-mapping

clang/examples/DPCT/Thrust/thrust$$max.cu

@@ -0,0 +1,21 @@
+#include <thrust/extrema.h>
+
+void max_test() {
+  // clang-format off
+  // Start
+  struct key_value {
+    int key;
+    int value;
+  };
+  struct compare_key_value {
+    __host__ __device__ bool operator()(key_value lhs, key_value rhs) {
+      return lhs.key < rhs.key;
+    }
+  };
+  key_value a = {13, 0};
+  key_value b = {7, 1};
+  key_value smaller = thrust::max(a, b, compare_key_value());
+  int value = thrust::max(1, 2);
+  // End
+  // clang-format on
+}

clang/examples/DPCT/Thrust/thrust$$min.cu

@@ -0,0 +1,21 @@
+#include <thrust/extrema.h>
+
+void min_test() {
+  // clang-format off
+  // Start
+  struct key_value {
+    int key;
+    int value;
+  };
+  struct compare_key_value {
+    __host__ __device__ bool operator()(key_value lhs, key_value rhs) {
+      return lhs.key < rhs.key;
+    }
+  };
+  key_value a = {13, 0};
+  key_value b = {7, 1};
+  key_value smaller = thrust::min(a, b, compare_key_value());
+  int value = thrust::min(1, 2);
+  // End
+  // clang-format on
+}

clang/lib/DPCT/RulesLangLib/APINamesThrust.inc

@@ -1172,10 +1172,18 @@ thrustFactory("thrust::detail::vector_equal",
               {{3,PolicyState::NoPolicy,3,"oneapi::dpl::equal", HelperFeatureEnum::none}}),
 
 // thrust::max
-CALL_FACTORY_ENTRY("thrust::max", CALL("std::max", ARG(0), ARG(1)))
+CONDITIONAL_FACTORY_ENTRY(
+   CheckArgCount(2), 
+   CALL_FACTORY_ENTRY("thrust::max", CALL("std::max", ARG(0), ARG(1))),
+   CALL_FACTORY_ENTRY("thrust::max", CALL("std::max", ARG(0), ARG(1), ARG(2)))
+)
 
 // thrust::min
-CALL_FACTORY_ENTRY("thrust::min", CALL("std::min",ARG(0), ARG(1)))
+CONDITIONAL_FACTORY_ENTRY(
+   CheckArgCount(2), 
+   CALL_FACTORY_ENTRY("thrust::min", CALL("std::min", ARG(0), ARG(1))),
+   CALL_FACTORY_ENTRY("thrust::min", CALL("std::min", ARG(0), ARG(1), ARG(2)))
+)
 
 // thrust::tie
 CALL_FACTORY_ENTRY("thrust::tie", CALL("std::tie",ARG(0), ARG(1)))

clang/test/dpct/query_api_mapping/Thrust/thrust_api_test_p3.cu

@@ -302,4 +302,62 @@
 // thrust_get_temporary_buffer-NEXT:   dpct::device_sys_tag device_sys;
 // thrust_get_temporary_buffer-NEXT:   ptr_and_size_t ptr_and_size = dpct::get_temporary_allocation<int>(device_sys, N);
 
-
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=thrust::max --extra-arg="-std=c++14"| FileCheck %s -check-prefix=thrust_max
+// thrust_max: CUDA API:
+// thrust_max-NEXT:   struct key_value {
+// thrust_max-NEXT:     int key;
+// thrust_max-NEXT:     int value;
+// thrust_max-NEXT:   };
+// thrust_max-NEXT:   struct compare_key_value {
+// thrust_max-NEXT:     __host__ __device__ bool operator()(key_value lhs, key_value rhs) {
+// thrust_max-NEXT:       return lhs.key < rhs.key;
+// thrust_max-NEXT:     }
+// thrust_max-NEXT:   };
+// thrust_max-NEXT:   key_value a = {13, 0};
+// thrust_max-NEXT:   key_value b = {7, 1};
+// thrust_max-NEXT:   key_value smaller = thrust::max(a, b, compare_key_value());
+// thrust_max-NEXT:   int value = thrust::max(1, 2);
+// thrust_max-NEXT: Is migrated to:
+// thrust_max-NEXT:   struct key_value {
+// thrust_max-NEXT:     int key;
+// thrust_max-NEXT:     int value;
+// thrust_max-NEXT:   };
+// thrust_max-NEXT:   struct compare_key_value {
+// thrust_max-NEXT:     bool operator()(key_value lhs, key_value rhs) {
+// thrust_max-NEXT:       return lhs.key < rhs.key;
+// thrust_max-NEXT:     }
+// thrust_max-NEXT:   };
+// thrust_max-NEXT:   key_value a = {13, 0};
+// thrust_max-NEXT:   key_value b = {7, 1};
+// thrust_max-NEXT:   key_value smaller = std::max(a, b, compare_key_value());
+// thrust_max-NEXT:   int value = std::max(1, 2);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=thrust::min --extra-arg="-std=c++14"| FileCheck %s -check-prefix=thrust_min
+// thrust_min: CUDA API:
+// thrust_min-NEXT:   struct key_value {
+// thrust_min-NEXT:     int key;
+// thrust_min-NEXT:     int value;
+// thrust_min-NEXT:   };
+// thrust_min-NEXT:   struct compare_key_value {
+// thrust_min-NEXT:     __host__ __device__ bool operator()(key_value lhs, key_value rhs) {
+// thrust_min-NEXT:       return lhs.key < rhs.key;
+// thrust_min-NEXT:     }
+// thrust_min-NEXT:   };
+// thrust_min-NEXT:   key_value a = {13, 0};
+// thrust_min-NEXT:   key_value b = {7, 1};
+// thrust_min-NEXT:   key_value smaller = thrust::min(a, b, compare_key_value());
+// thrust_min-NEXT:   int value = thrust::min(1, 2);
+// thrust_min-NEXT: Is migrated to:
+// thrust_min-NEXT:   struct key_value {
+// thrust_min-NEXT:     int key;
+// thrust_min-NEXT:     int value;
+// thrust_min-NEXT:   };
+// thrust_min-NEXT:   struct compare_key_value {
+// thrust_min-NEXT:     bool operator()(key_value lhs, key_value rhs) {
+// thrust_min-NEXT:       return lhs.key < rhs.key;
+// thrust_min-NEXT:     }
+// thrust_min-NEXT:   };
+// thrust_min-NEXT:   key_value a = {13, 0};
+// thrust_min-NEXT:   key_value b = {7, 1};
+// thrust_min-NEXT:   key_value smaller = std::min(a, b, compare_key_value());
+// thrust_min-NEXT:   int value = std::min(1, 2);

clang/test/dpct/query_api_mapping/test_all.cu

@@ -2474,9 +2474,11 @@
 // CHECK-NEXT: thrust::make_tuple
 // CHECK-NEXT: thrust::make_zip_iterator
 // CHECK-NEXT: thrust::malloc
+// CHECK-NEXT: thrust::max
 // CHECK-NEXT: thrust::max_element
 // CHECK-NEXT: thrust::merge
 // CHECK-NEXT: thrust::merge_by_key
+// CHECK-NEXT: thrust::min
 // CHECK-NEXT: thrust::min_element
 // CHECK-NEXT: thrust::minmax_element
 // CHECK-NEXT: thrust::mismatch

clang/test/dpct/thrust-for-h2o4gpu.cu

@@ -14,7 +14,7 @@
 #include <thrust/execution_policy.h>
 #include <thrust/random.h>
 #include <thrust/reduce.h>
-#include <algorithm>
+#include <thrust/extrema.h>
 #include <thrust/inner_product.h>
 #include <thrust/extrema.h>
 #include <thrust/host_vector.h>
@@ -651,3 +651,34 @@ template <class T> auto foo2(T t) {
   // CHECK: return dpct::make_constant_iterator(std::make_tuple(t, false));
   return thrust::make_constant_iterator(thrust::make_tuple<T, bool>(t, false));
 }
+
+struct key_value {
+   int key;
+   int value;
+};
+
+// CHECK: struct compare_key_value {
+// CHECK-NEXT:   bool operator()(key_value lhs, key_value rhs) {
+// CHECK-NEXT:     return lhs.key < rhs.key;
+// CHECK-NEXT:   }
+// CHECK-NEXT: };
+struct compare_key_value {
+  __host__ __device__ bool operator()(key_value lhs, key_value rhs) {
+    return lhs.key < rhs.key;
+  }
+};
+
+void thrust_max_min() {
+  key_value a = {13, 0};
+  key_value b = {7, 1};
+
+// CHECK:  key_value smaller = std::min(a, b, compare_key_value());
+// CHECK-NEXT:  key_value maxer = std::max(a, b, compare_key_value());
+  key_value smaller = thrust::min(a, b, compare_key_value());
+  key_value maxer = thrust::max(a, b, compare_key_value());
+
+// CHECK:   int min = std::min(1, 2);
+// CHECK-NEXT:  int max = std::max(1, 2);
+  int min = thrust::min(1, 2);
+  int max = thrust::max(1, 2);
+}