From 8bdc40b106896ed04012209db8024c9755747918 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 2 Sep 2020 11:12:31 -0400
Subject: [PATCH 1/4] [SYCL] Add is_native_function_object trait

These traits were mentioned in the GroupAlgorithms extension and the SYCL 2020
provisional specification, but weren't implemented anywhere.

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 .../CL/sycl/ONEAPI/group_algorithm.hpp        | 17 +++++++++++++++
 sycl/test/group-algorithm/traits.cpp          | 21 +++++++++++++++++++
 2 files changed, 38 insertions(+)
 create mode 100644 sycl/test/group-algorithm/traits.cpp

diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
index d93b068c09378..d6d7e9a39a137 100644
--- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
@@ -190,6 +190,23 @@ using EnableIfIsNonNativeOp = cl::sycl::detail::enable_if_t<
         !cl::sycl::detail::is_native_op<T, BinaryOperation>::value,
     T>;
 
+template <typename T> struct is_native_function_object : std::false_type {};
+
+template <template <typename> class BinaryOperation, typename T>
+struct is_native_function_object<BinaryOperation<T>> {
+  static const bool value =
+      (sycl::detail::is_scalar_arithmetic<T>::value ||
+       sycl::detail::is_vector_arithmetic<T>::value ||
+       std::is_same<T, void>::value) &&
+      sycl::detail::is_native_op<T, BinaryOperation<T>>::value;
+};
+
+#if __cplusplus >= 201703L
+template <typename T>
+inline constexpr bool is_native_function_object_v =
+    is_native_function_object<T>::value;
+#endif
+
 template <typename Group> bool all_of(Group, bool pred) {
   static_assert(sycl::detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
diff --git a/sycl/test/group-algorithm/traits.cpp b/sycl/test/group-algorithm/traits.cpp
new file mode 100644
index 0000000000000..0b1db4d591012
--- /dev/null
+++ b/sycl/test/group-algorithm/traits.cpp
@@ -0,0 +1,21 @@
+// RUN: %clangxx -std=c++17 -fsyntax-only -Xclang -verify %s -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning
+// expected-no-diagnostics
+#include <CL/sycl.hpp>
+
+using namespace sycl::ONEAPI;
+
+struct CustomType {};
+
+template <typename T> struct CustomFunctor {};
+
+int main() {
+  static_assert(is_native_function_object_v<plus<>>);
+  static_assert(is_native_function_object_v<multiplies<>>);
+  static_assert(is_native_function_object_v<bit_or<>>);
+  static_assert(is_native_function_object_v<bit_xor<>>);
+  static_assert(is_native_function_object_v<bit_and<>>);
+  static_assert(is_native_function_object_v<plus<float>>);
+  static_assert(is_native_function_object_v<plus<sycl::vec<float, 4>>>);
+  static_assert(!is_native_function_object_v<plus<CustomType>>);
+  static_assert(!is_native_function_object_v<CustomFunctor<float>>);
+}

From 81028155d0a7f3cfc8eda7e9a2692885f02c3b4a Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 2 Sep 2020 13:46:54 -0400
Subject: [PATCH 2/4] [SYCL] Use std::integral_constant

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
index d6d7e9a39a137..ed95ce9ca24ba 100644
--- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
@@ -193,12 +193,12 @@ using EnableIfIsNonNativeOp = cl::sycl::detail::enable_if_t<
 template <typename T> struct is_native_function_object : std::false_type {};
 
 template <template <typename> class BinaryOperation, typename T>
-struct is_native_function_object<BinaryOperation<T>> {
-  static const bool value =
-      (sycl::detail::is_scalar_arithmetic<T>::value ||
-       sycl::detail::is_vector_arithmetic<T>::value ||
-       std::is_same<T, void>::value) &&
-      sycl::detail::is_native_op<T, BinaryOperation<T>>::value;
+struct is_native_function_object<BinaryOperation<T>>
+    : std::integral_constant<
+          bool, (sycl::detail::is_scalar_arithmetic<T>::value ||
+                 sycl::detail::is_vector_arithmetic<T>::value ||
+                 std::is_same<T, void>::value) &&
+                    sycl::detail::is_native_op<T, BinaryOperation<T>>::value> {
 };
 
 #if __cplusplus >= 201703L

From 666bcbf0e31dfeaf17f607f0212631ea90fb5dab Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 2 Sep 2020 13:50:34 -0400
Subject: [PATCH 3/4] [SYCL] Add support for const via std::decay

Signed-off-by: John Pennycook <john.pennycook@intel.com>
Co-authored-by: Ruslan Arutyunyan <ruslan.arutyunyan@intel.com>
---
 sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp | 9 +++++++--
 sycl/test/group-algorithm/traits.cpp            | 1 +
 2 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
index ed95ce9ca24ba..c5ba16ec56ad7 100644
--- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
@@ -190,10 +190,11 @@ using EnableIfIsNonNativeOp = cl::sycl::detail::enable_if_t<
         !cl::sycl::detail::is_native_op<T, BinaryOperation>::value,
     T>;
 
-template <typename T> struct is_native_function_object : std::false_type {};
+template <typename T>
+struct is_native_function_object_impl : std::false_type {};
 
 template <template <typename> class BinaryOperation, typename T>
-struct is_native_function_object<BinaryOperation<T>>
+struct is_native_function_object_impl<BinaryOperation<T>>
     : std::integral_constant<
           bool, (sycl::detail::is_scalar_arithmetic<T>::value ||
                  sycl::detail::is_vector_arithmetic<T>::value ||
@@ -201,6 +202,10 @@ struct is_native_function_object<BinaryOperation<T>>
                     sycl::detail::is_native_op<T, BinaryOperation<T>>::value> {
 };
 
+template <typename T>
+struct is_native_function_object
+    : is_native_function_object_impl<std::decay_t<T>> {};
+
 #if __cplusplus >= 201703L
 template <typename T>
 inline constexpr bool is_native_function_object_v =
diff --git a/sycl/test/group-algorithm/traits.cpp b/sycl/test/group-algorithm/traits.cpp
index 0b1db4d591012..431b5bb22ab7c 100644
--- a/sycl/test/group-algorithm/traits.cpp
+++ b/sycl/test/group-algorithm/traits.cpp
@@ -14,6 +14,7 @@ int main() {
   static_assert(is_native_function_object_v<bit_or<>>);
   static_assert(is_native_function_object_v<bit_xor<>>);
   static_assert(is_native_function_object_v<bit_and<>>);
+  static_assert(is_native_function_object_v<const plus<>>);
   static_assert(is_native_function_object_v<plus<float>>);
   static_assert(is_native_function_object_v<plus<sycl::vec<float, 4>>>);
   static_assert(!is_native_function_object_v<plus<CustomType>>);

From 2037dbb48573cf25b56445af6825fc0c6839c2f9 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 2 Sep 2020 14:46:46 -0400
Subject: [PATCH 4/4] [SYCL] Fix C++11 include

std::decay_t not available until C++14

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
index c5ba16ec56ad7..c1203ea7044b7 100644
--- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
@@ -204,7 +204,7 @@ struct is_native_function_object_impl<BinaryOperation<T>>
 
 template <typename T>
 struct is_native_function_object
-    : is_native_function_object_impl<std::decay_t<T>> {};
+    : is_native_function_object_impl<typename std::decay<T>::type> {};
 
 #if __cplusplus >= 201703L
 template <typename T>