[SYCL] Refactor C++14 support

sycl/doc/PreprocessorMacros.md

@@ -31,7 +31,12 @@ SYCL 2020.
 
 ### `SYCL_DISABLE_DEPRECATION_WARNINGS`
 
-Disables all deprecation warnings in SYCL runtime headers, including SYCL 1.2.1 deprecations.
+Disables all deprecation warnings in SYCL runtime headers, including SYCL 1.2.1
+deprecations.
+
+### `SYCL_DISABLE_CPP_VERSION_CHECK_MESSAGE`
+
+Disables a message which warns about unsupported C++ version.
 
 ### Version macros
 

sycl/include/CL/sycl/detail/cg.hpp

@@ -79,8 +79,8 @@ namespace detail {
 //
 // 0x00000001 - CG type KERNEL version 0
 // 0x01000001 - CG type KERNEL version 1
-//   /\
-//   ||
+//    ^
+//    |
 // The byte specifies the version
 //
 // A user of this vector should not expect that a specific data is stored at a

sycl/include/CL/sycl/detail/cg_types.hpp

@@ -169,39 +169,43 @@ constexpr bool isKernelLambdaCallableWithKernelHandlerImpl() {
                                         kernel_handler>();
 }
 
-// Type traits to find out if kernal lambda has kernel_handler argument
+// Type trait to find out if kernal lambda has kernel_handler argument
+template <typename KernelType, typename LambdaArgType = void>
+struct KernelLambdaHasKernelHandlerArgT {
+  constexpr static bool value =
+      isKernelLambdaCallableWithKernelHandlerImpl<KernelType, LambdaArgType>();
+};
+
+// Helpers for running kernel lambda on the host device
 
 template <typename KernelType>
-constexpr bool isKernelLambdaCallableWithKernelHandler() {
-  return check_kernel_lambda_takes_args<KernelType, kernel_handler>();
+typename detail::enable_if_t<
+    KernelLambdaHasKernelHandlerArgT<KernelType>::value, void>
+runKernelWithoutArg(KernelType KernelName) {
+  kernel_handler KH;
+  KernelName(KH);
 }
 
-template <typename KernelType, typename LambdaArgType>
-constexpr bool isKernelLambdaCallableWithKernelHandler() {
-  return isKernelLambdaCallableWithKernelHandlerImpl<KernelType,
-                                                     LambdaArgType>();
+template <typename KernelType>
+typename detail::enable_if_t<
+    !KernelLambdaHasKernelHandlerArgT<KernelType>::value, void>
+runKernelWithoutArg(KernelType KernelName) {
+  KernelName();
 }
 
-// Helpers for running kernel lambda on the host device
-
-template <typename KernelType> void runKernelWithoutArg(KernelType KernelName) {
-  if constexpr (isKernelLambdaCallableWithKernelHandler<KernelType>()) {
-    kernel_handler KH;
-    KernelName(KH);
-  } else {
-    KernelName();
-  }
+template <typename ArgType, typename KernelType>
+typename detail::enable_if_t<
+    KernelLambdaHasKernelHandlerArgT<KernelType, ArgType>::value, void>
+runKernelWithArg(KernelType KernelName, ArgType Arg) {
+  kernel_handler KH;
+  KernelName(Arg, KH);
 }
 
 template <typename ArgType, typename KernelType>
-constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
-  if constexpr (isKernelLambdaCallableWithKernelHandler<KernelType,
-                                                        ArgType>()) {
-    kernel_handler KH;
-    KernelName(Arg, KH);
-  } else {
-    KernelName(Arg);
-  }
+typename detail::enable_if_t<
+    !KernelLambdaHasKernelHandlerArgT<KernelType, ArgType>::value, void>
+runKernelWithArg(KernelType KernelName, ArgType Arg) {
+  KernelName(Arg);
 }
 
 // The pure virtual class aimed to store lambda/functors of any type.

sycl/include/CL/sycl/detail/defines_elementary.hpp

@@ -39,7 +39,9 @@
 #endif
 
 #ifndef __SYCL_DEPRECATED
-#ifndef SYCL_DISABLE_DEPRECATION_WARNINGS
+// The deprecated attribute is not supported in some situations(e.g. namespace)
+// in C++14 mode
+#if !defined(SYCL2020_DISABLE_DEPRECATION_WARNINGS) && __cplusplus >= 201703L
 #define __SYCL_DEPRECATED(message) [[deprecated(message)]]
 #else // SYCL_DISABLE_DEPRECATION_WARNINGS
 #define __SYCL_DEPRECATED(message)
@@ -90,7 +92,7 @@
 // define __SYCL_WARNING convenience macro to report compiler warnings
 #if defined(__GNUC__)
 #define __SYCL_GCC_PRAGMA(x) _Pragma(#x)
-#define __SYCL_WARNING(msg) __SYCL_GCC_PRAGMA(GCC warning msg)
+#define __SYCL_WARNING(msg) __SYCL_GCC_PRAGMA(message msg)
 #elif defined(_MSC_VER) && !defined(__clang__)
 #define __SYCL_QUOTE1(x) #x
 #define __SYCL_QUOTE(x) __SYCL_QUOTE1(x)
@@ -100,3 +102,10 @@
 // clang emits "warning:" in the message pragma output
 #define __SYCL_WARNING(msg) __pragma(message(msg))
 #endif // __GNUC__
+
+// Emitting a message instead of a warning here in order to avoid hard error
+// if an app is compiled with "-werror" option.
+#if !defined(SYCL_DISABLE_CPP_VERSION_CHECK_MESSAGE) && __cplusplus < 201703L
+__SYCL_WARNING("DPCPP does not support C++ version earlier than C++17. Some "
+               "features might not be available.")
+#endif

sycl/include/CL/sycl/handler.hpp

@@ -524,8 +524,8 @@ class __SYCL_EXPORT handler {
   template <typename KernelName, typename KernelType, int Dims,
             typename LambdaArgType>
   void StoreLambda(KernelType KernelFunc) {
-    if (detail::isKernelLambdaCallableWithKernelHandler<KernelType,
-                                                        LambdaArgType>() &&
+    if (detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+                                                 LambdaArgType>::value &&
         MIsHost) {
       throw cl::sycl::feature_not_supported(
           "kernel_handler is not yet supported by host device.",
@@ -853,8 +853,8 @@ class __SYCL_EXPORT handler {
       setType(detail::CG::Kernel);
 #endif
     } else
-#endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ &&                     \
-       // !DPCPP_HOST_DEVICE_OPENMP && !DPCPP_HOST_DEVICE_PERF_NATIVE &&       \
+#endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ &&
+       // !DPCPP_HOST_DEVICE_OPENMP && !DPCPP_HOST_DEVICE_PERF_NATIVE &&
        // SYCL_LANGUAGE_VERSION >= 202001
     {
       (void)NumWorkItems;
@@ -1002,7 +1002,8 @@ class __SYCL_EXPORT handler {
   // Wrappers for kernel_single_task(...)
 
   template <typename KernelName, typename KernelType>
-  void
+  detail::enable_if_t<
+      detail::KernelLambdaHasKernelHandlerArgT<KernelType>::value>
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_single_task_wrapper(KernelType KernelFunc) {
 #else
@@ -1011,19 +1012,29 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     detail::CheckDeviceCopyable<KernelType>();
 #endif // __SYCL_DEVICE_ONLY__
-    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
-                      KernelType>()) {
-      kernel_handler KH;
-      kernel_single_task<KernelName>(KernelFunc, KH);
-    } else {
-      kernel_single_task<KernelName>(KernelFunc);
-    }
+    kernel_handler KH;
+    kernel_single_task<KernelName>(KernelFunc, KH);
+  }
+
+  template <typename KernelName, typename KernelType>
+  detail::enable_if_t<
+      !detail::KernelLambdaHasKernelHandlerArgT<KernelType>::value>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_single_task_wrapper(KernelType KernelFunc) {
+#else
+  kernel_single_task_wrapper(const KernelType &KernelFunc) {
+#endif
+#ifdef __SYCL_DEVICE_ONLY__
+    detail::CheckDeviceCopyable<KernelType>();
+#endif // __SYCL_DEVICE_ONLY__
+    kernel_single_task<KernelName>(KernelFunc);
   }
 
   // Wrappers for kernel_parallel_for(...)
 
   template <typename KernelName, typename ElementType, typename KernelType>
-  void
+  detail::enable_if_t<
+      detail::KernelLambdaHasKernelHandlerArgT<KernelType, ElementType>::value>
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_wrapper(KernelType KernelFunc) {
 #else
@@ -1032,19 +1043,29 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     detail::CheckDeviceCopyable<KernelType>();
 #endif // __SYCL_DEVICE_ONLY__
-    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
-                      KernelType, ElementType>()) {
-      kernel_handler KH;
-      kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
-    } else {
-      kernel_parallel_for<KernelName, ElementType>(KernelFunc);
-    }
+    kernel_handler KH;
+    kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
+  }
+
+  template <typename KernelName, typename ElementType, typename KernelType>
+  detail::enable_if_t<
+      !detail::KernelLambdaHasKernelHandlerArgT<KernelType, ElementType>::value>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_wrapper(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
+#endif
+#ifdef __SYCL_DEVICE_ONLY__
+    detail::CheckDeviceCopyable<KernelType>();
+#endif // __SYCL_DEVICE_ONLY__
+    kernel_parallel_for<KernelName, ElementType>(KernelFunc);
   }
 
   // Wrappers for kernel_parallel_for_work_group(...)
 
   template <typename KernelName, typename ElementType, typename KernelType>
-  void
+  detail::enable_if_t<
+      detail::KernelLambdaHasKernelHandlerArgT<KernelType, ElementType>::value>
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
 #else
@@ -1053,13 +1074,22 @@ class __SYCL_EXPORT handler {
 #ifdef __SYCL_DEVICE_ONLY__
     detail::CheckDeviceCopyable<KernelType>();
 #endif // __SYCL_DEVICE_ONLY__
-    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
-                      KernelType, ElementType>()) {
-      kernel_handler KH;
-      kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc, KH);
-    } else {
-      kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
-    }
+    kernel_handler KH;
+    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc, KH);
+  }
+
+  template <typename KernelName, typename ElementType, typename KernelType>
+  detail::enable_if_t<
+      !detail::KernelLambdaHasKernelHandlerArgT<KernelType, ElementType>::value>
+#ifdef __SYCL_NONCONST_FUNCTOR__
+  kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
+#else
+  kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
+#endif
+#ifdef __SYCL_DEVICE_ONLY__
+    detail::CheckDeviceCopyable<KernelType>();
+#endif // __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
   }
 
   std::shared_ptr<detail::kernel_bundle_impl>
@@ -2445,25 +2475,31 @@ class __SYCL_EXPORT handler {
   friend class ::MockHandler;
 
   template <typename WrapperT, typename TransformedArgType, int Dims,
-            typename KernelType>
+            typename KernelType,
+            detail::enable_if_t<detail::KernelLambdaHasKernelHandlerArgT<
+                KernelType, TransformedArgType>::value> * = nullptr>
   auto getRangeRoundedKernelLambda(KernelType KernelFunc,
                                    range<Dims> NumWorkItems) {
-    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
-                      KernelType, TransformedArgType>()) {
-      return [=](TransformedArgType Arg, kernel_handler KH) {
-        if (Arg[0] >= NumWorkItems[0])
-          return;
-        Arg.set_allowed_range(NumWorkItems);
-        KernelFunc(Arg, KH);
-      };
-    } else {
-      return [=](TransformedArgType Arg) {
-        if (Arg[0] >= NumWorkItems[0])
-          return;
-        Arg.set_allowed_range(NumWorkItems);
-        KernelFunc(Arg);
-      };
-    }
+    return [=](TransformedArgType Arg, kernel_handler KH) {
+      if (Arg[0] >= NumWorkItems[0])
+        return;
+      Arg.set_allowed_range(NumWorkItems);
+      KernelFunc(Arg, KH);
+    };
+  }
+
+  template <typename WrapperT, typename TransformedArgType, int Dims,
+            typename KernelType,
+            detail::enable_if_t<!detail::KernelLambdaHasKernelHandlerArgT<
+                KernelType, TransformedArgType>::value> * = nullptr>
+  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
+                                   range<Dims> NumWorkItems) {
+    return [=](TransformedArgType Arg) {
+      if (Arg[0] >= NumWorkItems[0])
+        return;
+      Arg.set_allowed_range(NumWorkItems);
+      KernelFunc(Arg);
+    };
   }
 };
 } // namespace sycl

sycl/include/CL/sycl/id.hpp

@@ -109,10 +109,6 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
 
 // OP is: ==, !=
 #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
-  using detail::array<dimensions>::operator==;
-#if __cpp_impl_three_way_comparison < 201907
-  using detail::array<dimensions>::operator!=;
-#endif
 
   /* Enable operators with integral types.
    * Template operators take precedence than type conversion. In the case of
@@ -126,6 +122,11 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
       return false op true;                                                    \
     return true op true;                                                       \
   }                                                                            \
+  bool operator op(const id<dimensions> &rhs) const {                          \
+    if (this->common_array[0] != rhs)                                          \
+      return false op true;                                                    \
+    return true op true;                                                       \
+  }                                                                            \
   template <typename T>                                                        \
   friend EnableIfIntegral<T, bool> operator op(const T &lhs,                   \
                                                const id<dimensions> &rhs) {    \

sycl/test/basic_tests/stdcpp_compat.cpp

@@ -1,7 +1,7 @@
-// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s -c -o %t.out
-// RUN: %clangxx -fsycl -std=c++14 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s -c -o %t.out
-// RUN: %clangxx -fsycl -std=c++17 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s -c -o %t.out
-// RUN: %clangxx -fsycl -std=c++20 -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s -c -o %t.out
+// RUN: %clangxx -fsycl --no-system-header-prefix=CL/sycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify %s -c -o %t.out
+// RUN: %clangxx -std=c++14 -fsycl --no-system-header-prefix=CL/sycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify %s -c -o %t.out
+// RUN: %clangxx -std=c++17 -fsycl --no-system-header-prefix=CL/sycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify %s -c -o %t.out
+// RUN: %clangxx -std=c++20 -fsycl --no-system-header-prefix=CL/sycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify %s -c -o %t.out
 // expected-no-diagnostics
 
 #include <CL/sycl.hpp>