[SYCL] Make spec_constant default constructor public and available on host. (#2488)

This is what the implemented spec prescribes:
https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp

@@ -32,17 +32,14 @@ class spec_const_error : public compile_program_error {
 };
 
 template <typename T, typename ID = T> class spec_constant {
-private:
-  // Implementation defined constructor.
-#ifdef __SYCL_DEVICE_ONLY__
 public:
   spec_constant() {}
 
 private:
-#else
-  spec_constant(T Cst) : Val(Cst) {}
-#endif
 #ifndef __SYCL_DEVICE_ONLY__
+  // Implementation defined constructor.
+  spec_constant(T Cst) : Val(Cst) {}
+
   T Val;
 #endif
   friend class cl::sycl::program;

sycl/test/spec_const/spec_const_hw.cpp

@@ -23,16 +23,17 @@
 
 class MyInt32Const;
 class MyFloatConst;
+class MyConst;
 
 using namespace sycl;
 
 class KernelAAAi;
 class KernelBBBf;
 
-int val = 10;
+int global_val = 10;
 
 // Fetch a value at runtime.
-int get_value() { return val; }
+int get_value() { return global_val; }
 
 float foo(
     const cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst>
@@ -49,8 +50,22 @@ struct SCWrapper {
   cl::sycl::ONEAPI::experimental::spec_constant<int, class sc_name2> SC2;
 };
 
+// MyKernel is used to test default constructor
+using AccT = sycl::accessor<int, 1, sycl::access::mode::write>;
+using ScT = sycl::ONEAPI::experimental::spec_constant<int, MyConst>;
+
+struct MyKernel {
+  MyKernel(AccT &Acc) : Acc(Acc) {}
+
+  void setConst(ScT Sc) { this->Sc = Sc; }
+
+  void operator()() const { Acc[0] = Sc.get(); }
+  AccT Acc;
+  ScT Sc;
+};
+
 int main(int argc, char **argv) {
-  val = argc + 16;
+  global_val = argc + 16;
 
   cl::sycl::queue q(default_selector{}, [](exception_list l) {
     for (auto ep : l) {
@@ -68,10 +83,11 @@ int main(int argc, char **argv) {
 
   std::cout << "Running on " << q.get_device().get_info<info::device::name>()
             << "\n";
-  std::cout << "val = " << val << "\n";
+  std::cout << "global_val = " << global_val << "\n";
   cl::sycl::program program1(q.get_context());
   cl::sycl::program program2(q.get_context());
   cl::sycl::program program3(q.get_context());
+  cl::sycl::program program4(q.get_context());
 
   int goldi = (int)get_value();
   // TODO make this floating point once supported by the compiler
@@ -83,22 +99,30 @@ int main(int argc, char **argv) {
   cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst> f32 =
       program2.set_spec_constant<MyFloatConst>(goldf);
 
+  cl::sycl::ONEAPI::experimental::spec_constant<int, MyConst> sc =
+      program4.set_spec_constant<MyConst>(goldi);
+
   program1.build_with_kernel_type<KernelAAAi>();
   // Use an option (does not matter which exactly) to test different internal
   // SYCL RT execution path
   program2.build_with_kernel_type<KernelBBBf>("-cl-fast-relaxed-math");
 
   SCWrapper W(program3);
   program3.build_with_kernel_type<class KernelWrappedSC>();
+
+  program4.build_with_kernel_type<MyKernel>();
+
   int goldw = 6;
 
   std::vector<int> veci(1);
   std::vector<float> vecf(1);
   std::vector<int> vecw(1);
+  std::vector<int> vec(1);
   try {
     cl::sycl::buffer<int, 1> bufi(veci.data(), veci.size());
     cl::sycl::buffer<float, 1> buff(vecf.data(), vecf.size());
     cl::sycl::buffer<int, 1> bufw(vecw.data(), vecw.size());
+    cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());
 
     q.submit([&](cl::sycl::handler &cgh) {
       auto acci = bufi.get_access<cl::sycl::access::mode::write>(cgh);
@@ -123,6 +147,19 @@ int main(int argc, char **argv) {
           program3.get_kernel<KernelWrappedSC>(),
           [=]() { accw[0] = W.SC1.get() + W.SC2.get(); });
     });
+    // Check spec_constant default construction with subsequent initialization
+    q.submit([&](cl::sycl::handler &cgh) {
+      auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
+      // Specialization constants specification says:
+      //   cl::sycl::experimental::spec_constant is default constructible,
+      //   although the object is not considered initialized until the result of
+      //   the call to cl::sycl::program::set_spec_constant is assigned to it.
+      MyKernel Kernel(acc); // default construct inside MyKernel instance
+      Kernel.setConst(sc);  // initialize to sc, returned by set_spec_constant
+
+      cgh.single_task<MyKernel>(program4.get_kernel<MyKernel>(), Kernel);
+    });
+
   } catch (cl::sycl::exception &e) {
     std::cout << "*** Exception caught: " << e.what() << "\n";
     return 1;
@@ -146,6 +183,12 @@ int main(int argc, char **argv) {
     std::cout << "*** ERROR: " << valw << " != " << goldw << "(gold)\n";
     passed = false;
   }
+  int val = vec[0];
+
+  if (val != goldi) {
+    std::cout << "*** ERROR: " << val << " != " << goldi << "(gold)\n";
+    passed = false;
+  }
   std::cout << (passed ? "passed\n" : "FAILED\n");
   return passed ? 0 : 1;
 }