[SYCL] Add DPC++ RT support for SYCL 2020 spec constants (part 1)

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

@@ -149,23 +149,39 @@ static constexpr bool check_kernel_lambda_takes_args() {
   return check_fn_signature<std::remove_reference_t<F>, void(Args...)>::value;
 }
 
-// Type traits to find out if kernal lambda has kernel_handler argument
+// isKernelLambdaCallableWithKernelHandlerImpl checks if LambdaArgType is void
+// (e.g., in single_task), and based on that, calls
+// check_kernel_lambda_takes_args with proper set of arguments. Also this type
+// trait workarounds compilation error which happens only with msvc.
 
-template <typename KernelType, typename LambdaArgType = void,
+template <typename KernelType, typename LambdaArgType,
           typename std::enable_if_t<std::is_same<LambdaArgType, void>::value>
               * = nullptr>
-constexpr bool isKernelLambdaCallableWithKernelHandler() {
+constexpr bool isKernelLambdaCallableWithKernelHandlerImpl() {
   return check_kernel_lambda_takes_args<KernelType, kernel_handler>();
 }
 
 template <typename KernelType, typename LambdaArgType,
           typename std::enable_if_t<!std::is_same<LambdaArgType, void>::value>
               * = nullptr>
-constexpr bool isKernelLambdaCallableWithKernelHandler() {
+constexpr bool isKernelLambdaCallableWithKernelHandlerImpl() {
   return check_kernel_lambda_takes_args<KernelType, LambdaArgType,
                                         kernel_handler>();
 }
 
+// Type traits to find out if kernal lambda has kernel_handler argument
+
+template <typename KernelType>
+constexpr bool isKernelLambdaCallableWithKernelHandler() {
+  return check_kernel_lambda_takes_args<KernelType, kernel_handler>();
+}
+
+template <typename KernelType, typename LambdaArgType>
+constexpr bool isKernelLambdaCallableWithKernelHandler() {
+  return isKernelLambdaCallableWithKernelHandlerImpl<KernelType,
+                                                     LambdaArgType>();
+}
+
 // Helpers for running kernel lambda on the host device
 
 template <typename KernelType,

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -26,7 +26,8 @@ enum class kernel_param_kind_t {
   kind_accessor = 0,
   kind_std_layout = 1, // standard layout object parameters
   kind_sampler = 2,
-  kind_pointer = 3
+  kind_pointer = 3,
+  kind_specialization_constants_buffer = 4,
 };
 
 // describes a kernel parameter

sycl/include/CL/sycl/handler.hpp

@@ -498,7 +498,7 @@ class __SYCL_EXPORT handler {
                                                         LambdaArgType>() &&
         MIsHost) {
       throw cl::sycl::feature_not_supported(
-          "kernel_handler is not supported by host device.",
+          "kernel_handler is not yet supported by host device.",
           PI_INVALID_OPERATION);
     }
     MHostKernel.reset(

sycl/include/CL/sycl/kernel_handler.hpp

@@ -32,8 +32,9 @@ class kernel_handler {
 #ifdef __SYCL_DEVICE_ONLY__
     return getSpecializationConstantOnDevice<S>();
 #else
+    // TODO: add support of host device
     throw cl::sycl::feature_not_supported(
-        "kernel_handler::get_specialization_constant() is not supported by "
+        "kernel_handler::get_specialization_constant() is not yet supported by "
         "host device.",
         PI_INVALID_OPERATION);
 #endif // __SYCL_DEVICE_ONLY__

sycl/include/CL/sycl/specialization_id.hpp

@@ -20,15 +20,15 @@ template <typename T> class specialization_id {
 
   template <class... Args>
   explicit constexpr specialization_id(Args &&... args)
-      : MSpecializationConstantValue(args...) {}
+      : MDefaultValue(std::forward<Args>(args)...) {}
 
   specialization_id(const specialization_id &rhs) = delete;
   specialization_id(specialization_id &&rhs) = delete;
   specialization_id &operator=(const specialization_id &rhs) = delete;
   specialization_id &operator=(specialization_id &&rhs) = delete;
 
 private:
-  T MSpecializationConstantValue;
+  T MDefaultValue;
 };
 
 } // namespace sycl

sycl/source/detail/scheduler/commands.cpp

@@ -1695,6 +1695,12 @@ pi_result ExecCGCommand::SetKernelParamsAndLaunch(
                                                        Arg.MSize, Arg.MPtr);
       break;
     }
+    case kernel_param_kind_t::kind_specialization_constants_buffer: {
+      throw cl::sycl::feature_not_supported(
+        "Specialization constants are not yet fully supported",
+        PI_INVALID_OPERATION);
+      break;
+    }
     }
     ++NextTrueIndex;
   }

sycl/source/handler.cpp

@@ -236,6 +236,12 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
                        Index + IndexShift);
     break;
   }
+  case kernel_param_kind_t::kind_specialization_constants_buffer: {
+    throw cl::sycl::feature_not_supported(
+        "Specialization constants are not yet fully supported",
+        PI_INVALID_OPERATION);
+    break;
+  }
   }
 }
 

sycl/test/on-device/basic_tests/specialization_constants/kernel_lambda_with_kernel_handler_arg.cpp

@@ -1,6 +1,10 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %t.out
 
+// This test checks all possible scenarios of running single_task, parallel_for
+// and parallel_for_work_group to verify that this code compiles and runs
+// correctly with user's lambda with and without sycl::kernel_handler argument
+
 #include <CL/sycl.hpp>
 
 int main() {