[SYCL] Use if constexpr instead of enable_if to avoid msvc bug

Author: romanovvlad

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

@@ -184,34 +184,24 @@ constexpr bool isKernelLambdaCallableWithKernelHandler() {
 
 // Helpers for running kernel lambda on the host device
 
-template <typename KernelType,
-          typename std::enable_if_t<isKernelLambdaCallableWithKernelHandler<
-              KernelType>()> * = nullptr>
-constexpr void runKernelWithoutArg(KernelType KernelName) {
-  kernel_handler KH;
-  KernelName(KH);
-}
-
-template <typename KernelType,
-          typename std::enable_if_t<!isKernelLambdaCallableWithKernelHandler<
-              KernelType>()> * = nullptr>
-constexpr void runKernelWithoutArg(KernelType KernelName) {
-  KernelName();
-}
-
-template <typename ArgType, typename KernelType,
-          typename std::enable_if_t<isKernelLambdaCallableWithKernelHandler<
-              KernelType, ArgType>()> * = nullptr>
-constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
-  kernel_handler KH;
-  KernelName(Arg, KH);
+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<!isKernelLambdaCallableWithKernelHandler<
-              KernelType, ArgType>()> * = nullptr>
+template <typename ArgType, typename KernelType>
 constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
-  KernelName(Arg);
+  if constexpr (isKernelLambdaCallableWithKernelHandler<KernelType,
+                                                        ArgType>()) {
+    kernel_handler KH;
+    KernelName(Arg, KH);
+  } else {
+    KernelName(Arg);
+  }
 }
 
 // The pure virtual class aimed to store lambda/functors of any type.

sycl/include/CL/sycl/handler.hpp

@@ -968,80 +968,57 @@ class __SYCL_EXPORT handler {
   // Wrappers for kernel_single_task(...)
 
   template <typename KernelName, typename KernelType>
-  std::enable_if_t<
-      detail::isKernelLambdaCallableWithKernelHandler<KernelType>(), void>
-#ifdef __SYCL_NONCONST_FUNCTOR__
-  kernel_single_task_wrapper(KernelType KernelFunc) {
-#else
-  kernel_single_task_wrapper(const KernelType &KernelFunc) {
-#endif
-    kernel_handler KH;
-    kernel_single_task<KernelName>(KernelFunc, KH);
-  }
-
-  template <typename KernelName, typename KernelType>
-  std::enable_if_t<
-      !detail::isKernelLambdaCallableWithKernelHandler<KernelType>(), void>
+  void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_single_task_wrapper(KernelType KernelFunc) {
 #else
   kernel_single_task_wrapper(const KernelType &KernelFunc) {
 #endif
-    kernel_single_task<KernelName>(KernelFunc);
+    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
+                      KernelType>()) {
+      kernel_handler KH;
+      kernel_single_task<KernelName>(KernelFunc, KH);
+    } else {
+      kernel_single_task<KernelName>(KernelFunc);
+    }
   }
 
   // Wrappers for kernel_parallel_for(...)
 
   template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
-#ifdef __SYCL_NONCONST_FUNCTOR__
-  kernel_parallel_for_wrapper(KernelType KernelFunc) {
-#else
-  kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
-#endif
-    kernel_handler KH;
-    kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
-  }
-
-  template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<!detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
+  void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_wrapper(KernelType KernelFunc) {
 #else
   kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
 #endif
-    kernel_parallel_for<KernelName, ElementType>(KernelFunc);
+    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
+                      KernelType, ElementType>()) {
+      kernel_handler KH;
+      kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
+    }
+    else {
+      kernel_parallel_for<KernelName, ElementType>(KernelFunc);
+    }
   }
 
   // Wrappers for kernel_parallel_for_work_group(...)
 
   template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
-#ifdef __SYCL_NONCONST_FUNCTOR__
-  kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
-#else
-  kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
-#endif
-    kernel_handler KH;
-    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc, KH);
-  }
-
-  template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<!detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
+  void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
 #else
   kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
 #endif
-    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
+    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);
+    }
   }
 
   std::shared_ptr<detail::kernel_bundle_impl>
@@ -2289,32 +2266,25 @@ class __SYCL_EXPORT handler {
 
   friend class ::MockHandler;
 
-  template <
-      typename TransformedArgType, int Dims, typename KernelType,
-      typename std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
-          KernelType, TransformedArgType>()> * = nullptr>
+  template <typename TransformedArgType, int Dims, typename KernelType>
   auto getRangeRoundedKernelLambda(KernelType KernelFunc,
                                    range<Dims> NumWorkItems) {
-    return [=](TransformedArgType Arg, kernel_handler KH) {
-      if (Arg[0] >= NumWorkItems[0])
-        return;
-      Arg.set_allowed_range(NumWorkItems);
-      KernelFunc(Arg, KH);
-    };
-  }
-
-  template <typename TransformedArgType, int Dims, typename KernelType,
-            typename std::enable_if_t<
-                !detail::isKernelLambdaCallableWithKernelHandler<
-                    KernelType, TransformedArgType>()> * = nullptr>
-  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
-                                   range<Dims> NumWorkItems) {
-    return [=](TransformedArgType Arg) {
-      if (Arg[0] >= NumWorkItems[0])
-        return;
-      Arg.set_allowed_range(NumWorkItems);
-      KernelFunc(Arg);
-    };
+    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);
+      };
+    }
   }
 };
 } // namespace sycl