diff --git a/clang/examples/DPCT/Thrust/thrust$$max.cu b/clang/examples/DPCT/Thrust/thrust$$max.cu new file mode 100644 index 000000000000..8f19988919b1 --- /dev/null +++ b/clang/examples/DPCT/Thrust/thrust$$max.cu @@ -0,0 +1,21 @@ +#include + +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 +} diff --git a/clang/examples/DPCT/Thrust/thrust$$min.cu b/clang/examples/DPCT/Thrust/thrust$$min.cu new file mode 100644 index 000000000000..74bc063de6ed --- /dev/null +++ b/clang/examples/DPCT/Thrust/thrust$$min.cu @@ -0,0 +1,21 @@ +#include + +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 +} diff --git a/clang/lib/DPCT/RulesLangLib/APINamesThrust.inc b/clang/lib/DPCT/RulesLangLib/APINamesThrust.inc index 0778bbc5d4f4..8d50422848d0 100644 --- a/clang/lib/DPCT/RulesLangLib/APINamesThrust.inc +++ b/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))) diff --git a/clang/test/dpct/query_api_mapping/Thrust/thrust_api_test_p3.cu b/clang/test/dpct/query_api_mapping/Thrust/thrust_api_test_p3.cu index c033c2e1480c..48e38bee7a6f 100644 --- a/clang/test/dpct/query_api_mapping/Thrust/thrust_api_test_p3.cu +++ b/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(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); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index 25e69d6caef6..84da6a294711 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/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 diff --git a/clang/test/dpct/thrust-for-h2o4gpu.cu b/clang/test/dpct/thrust-for-h2o4gpu.cu index 8811847a021c..cf8423513a3b 100644 --- a/clang/test/dpct/thrust-for-h2o4gpu.cu +++ b/clang/test/dpct/thrust-for-h2o4gpu.cu @@ -14,7 +14,7 @@ #include #include #include -#include +#include #include #include #include @@ -651,3 +651,34 @@ template auto foo2(T t) { // CHECK: return dpct::make_constant_iterator(std::make_tuple(t, false)); return thrust::make_constant_iterator(thrust::make_tuple(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); +}