[SYCL] Refactor C++14 support

sycl/doc/PreprocessorMacros.md

@@ -31,7 +31,12 @@ This file describes macros that have effect on SYCL compiler and run-time.
 
 - **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_WARNING**
+
+  Disables a message which warns about unsupported C++ version.
 
 - **SYCL_DISABLE_FALLBACK_ASSERT**
 

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,41 @@ 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 std::enable_if_t<KernelLambdaHasKernelHandlerArgT<KernelType>::value>
+runKernelWithoutArg(KernelType KernelName) {
+  kernel_handler KH;
+  KernelName(KH);
 }
 
-template <typename KernelType, typename LambdaArgType>
-constexpr bool isKernelLambdaCallableWithKernelHandler() {
-  return isKernelLambdaCallableWithKernelHandlerImpl<KernelType,
-                                                     LambdaArgType>();
+template <typename KernelType>
+typename std::enable_if_t<!KernelLambdaHasKernelHandlerArgT<KernelType>::value>
+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 std::enable_if_t<
+    KernelLambdaHasKernelHandlerArgT<KernelType, ArgType>::value>
+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 std::enable_if_t<
+    !KernelLambdaHasKernelHandlerArgT<KernelType, ArgType>::value>
+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)
@@ -118,3 +120,17 @@
 #define __SYCL_UNROLL(x)
 #endif // compiler switch
 #endif // __SYCL_UNROLL
+
+#if !defined(SYCL_DISABLE_CPP_VERSION_CHECK_WARNING) && __cplusplus < 201703L
+
+#if defined(_MSC_VER) && !defined(__clang__)
+__SYCL_WARNING("DPCPP does not support C++ version earlier than C++17. Some "
+               "features might not be available.")
+#else
+// This is the only way to emit a warning from system headers using clang, it
+// cannot be wrapped by a macro(__pragma warning doesn't work in system
+// headers). The solution is borrowed from libcxx.
+#warning: DPCPP does not support C++ version earlier than C++17. Some features might not be available.
+#endif
+
+#endif

sycl/include/CL/sycl/handler.hpp

@@ -561,8 +561,9 @@ class __SYCL_EXPORT handler {
             typename LambdaArgType>
   void StoreLambda(KernelType KernelFunc) {
     constexpr bool IsCallableWithKernelHandler =
-        detail::isKernelLambdaCallableWithKernelHandler<KernelType,
-                                                        LambdaArgType>();
+        detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+                                                 LambdaArgType>::value;
+
     if (IsCallableWithKernelHandler && MIsHost) {
       throw cl::sycl::feature_not_supported(
           "kernel_handler is not yet supported by host device.",
@@ -591,7 +592,7 @@ class __SYCL_EXPORT handler {
 
     // If the kernel lambda is callable with a kernel_handler argument, manifest
     // the associated kernel handler.
-    if constexpr (IsCallableWithKernelHandler) {
+    if (IsCallableWithKernelHandler) {
       getOrInsertHandlerKernelBundle(/*Insert=*/true);
     }
   }
@@ -878,8 +879,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;
@@ -1027,7 +1028,7 @@ class __SYCL_EXPORT handler {
   // Wrappers for kernel_single_task(...)
 
   template <typename KernelName, typename KernelType>
-  void
+  std::enable_if_t<detail::KernelLambdaHasKernelHandlerArgT<KernelType>::value>
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_single_task_wrapper(KernelType KernelFunc) {
 #else
@@ -1036,19 +1037,28 @@ 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>
+  std::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
+  std::enable_if_t<
+      detail::KernelLambdaHasKernelHandlerArgT<KernelType, ElementType>::value>
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_wrapper(KernelType KernelFunc) {
 #else
@@ -1057,19 +1067,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>
+  std::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
+  std::enable_if_t<
+      detail::KernelLambdaHasKernelHandlerArgT<KernelType, ElementType>::value>
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
 #else
@@ -1078,13 +1098,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>
+  std::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>
@@ -2486,18 +2515,24 @@ class __SYCL_EXPORT handler {
                                 size_t &MinRange);
 
   template <typename WrapperT, typename TransformedArgType, int Dims,
-            typename KernelType>
+            typename KernelType,
+            std::enable_if_t<detail::KernelLambdaHasKernelHandlerArgT<
+                KernelType, TransformedArgType>::value> * = nullptr>
   auto getRangeRoundedKernelLambda(KernelType KernelFunc,
                                    range<Dims> NumWorkItems) {
-    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
-                      KernelType, TransformedArgType>()) {
-      return detail::RoundedRangeKernelWithKH<TransformedArgType, Dims,
-                                              KernelType>(NumWorkItems,
-                                                          KernelFunc);
-    } else {
-      return detail::RoundedRangeKernel<TransformedArgType, Dims, KernelType>(
-          NumWorkItems, KernelFunc);
-    }
+    return detail::RoundedRangeKernelWithKH<TransformedArgType, Dims,
+                                            KernelType>(NumWorkItems,
+                                                        KernelFunc);
+  }
+
+  template <typename WrapperT, typename TransformedArgType, int Dims,
+            typename KernelType,
+            std::enable_if_t<!detail::KernelLambdaHasKernelHandlerArgT<
+                KernelType, TransformedArgType>::value> * = nullptr>
+  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
+                                   range<Dims> NumWorkItems) {
+    return detail::RoundedRangeKernel<TransformedArgType, Dims, KernelType>(
+        NumWorkItems, KernelFunc);
   }
 };
 } // namespace sycl

sycl/test/basic_tests/stdcpp_compat.cpp

@@ -1,7 +1,38 @@
-// 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
-// expected-no-diagnostics
+// RUN: %clangxx -std=c++14 -fsycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify=expected,cxx14 %s -c -o %t.out
+// RUN: %clangxx -std=c++14 -fsycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify %s -c -o %t.out -DSYCL_DISABLE_CPP_VERSION_CHECK_WARNING=1
+// RUN: %clangxx -std=c++14 -fsycl --no-system-header-prefix=CL/sycl -Wall -pedantic -Wno-c99-extensions -Wno-deprecated -fsyntax-only -Xclang -verify=cxx14,warning_extension,expected %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
+// 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
+
+// The test checks SYCL headers C++ compiance and that a warning is emitted
+// when compiling in < C++17 mode.
 
 #include <CL/sycl.hpp>
+
+// cxx14-warning@* {{DPCPP does not support C++ version earlier than C++17. Some features might not be available.}}
+//
+// The next warning is not emitted in device compilation for some reason
+// warning_extension-warning@* 0-1 {{#warning is a language extension}}
+//
+// The next warning is emitted for windows only
+// expected-warning@* 0-1 {{Alignment of class vec is not in accordance with SYCL specification requirements, a limitation of the MSVC compiler(Error C2719).Applied default alignment.}}
+
+
+class KernelName1;
+
+int main() {
+
+  // Simple code to trigger instantiation of basic classes
+  sycl::buffer<int, 1> Buf(sycl::range<1>{42});
+
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &CGH) {
+    auto Acc = Buf.get_access<cl::sycl::access::mode::read>(CGH);
+
+    CGH.parallel_for<KernelName1>(sycl::range<1>{42},
+                                  [=](sycl::id<1> ID) { (void)Acc; });
+  });
+  Q.wait();
+}