[SYCL] Add specialization constant support in SYCL runtime.

Based on
https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md

1. Define SYCL API (sycl/include/CL/sycl/experimental/spec_constant.hpp)
2. Add convenience C++ wrappers for PI device binary structures and refactor
   runtime to use the wrappers. Get rid of custom deleters for binary images.
3. Implement SYCL spec constant APIs in program an program manager.

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

Author: kbobrovs

sycl/CMakeLists.txt

@@ -210,7 +210,7 @@ add_custom_target( sycl-toolchain
           clang
           clang-offload-wrapper
           clang-offload-bundler
-	  file-table-tform
+          file-table-tform
           llc
           llvm-ar
           llvm-foreach

sycl/include/CL/sycl/detail/os_util.hpp

@@ -80,6 +80,10 @@ class OSUtil {
   /// single one at most.
   static constexpr OSModuleHandle ExeModuleHandle = -1;
 
+  /// Dummy module handle to designate non-existing module for a device binary
+  /// image loaded from file e.g. via SYCL_USE_KERNEL_SPV env var.
+  static constexpr OSModuleHandle DummyModuleHandle = -2;
+
 #ifdef SYCL_RT_OS_WINDOWS
   static constexpr const char* DirSep = "\\";
 #else

sycl/include/CL/sycl/detail/spec_constant_impl.hpp

@@ -0,0 +1,46 @@
+//==-- spec_constant_impl.hpp - SYCL RT model for specialization constants -==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/detail/defines.hpp>
+
+#include <iostream>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// Represents a specialization constant in SYCL runtime.
+class spec_constant_impl {
+public:
+  spec_constant_impl(unsigned int ID) : ID(ID), Size(0), Bytes{0} {}
+
+  spec_constant_impl(unsigned int ID, size_t Size, const void *Val) : ID(ID) {
+    set(Size, Val);
+  }
+
+  void set(size_t Size, const void *Val);
+
+  unsigned int getID() const { return ID; }
+  size_t getSize() const { return Size; }
+  const unsigned char *getValuePtr() const { return Bytes; }
+  bool isSet() const { return Size != 0; }
+
+private:
+  unsigned int ID; // specialization constant's ID (equals to SPIRV ID)
+  size_t Size;     // size of its value
+  // TODO invent more flexible approach to support values of arbitrary type:
+  unsigned char Bytes[8]; // memory to hold the value bytes
+};
+
+std::ostream &operator<<(std::ostream &Out, const spec_constant_impl &V);
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp

@@ -12,13 +12,10 @@
 
 #ifdef __SYCL_DEVICE_ONLY__
 
-// Returns a unique string identifying the template parameter type. Stable
-// across device compiler invocations.
-template <typename T> const char *__sycl_fe_getStableUniqueTypeName();
-
 // Get the value of the specialization constant with given name.
 // Post-link tool traces the ID to a string literal it points to and assigns
 // integer ID.
-template <typename T> T __sycl_getSpecConstantValue(const char *ID);
+template <typename T>
+SYCL_EXTERNAL T __sycl_getSpecConstantValue(const char *ID);
 
 #endif

sycl/include/CL/sycl/detail/util.hpp

@@ -12,6 +12,7 @@
 
 #include <CL/sycl/detail/defines.hpp>
 
+#include <cstring>
 #include <mutex>
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -30,6 +31,34 @@ class Sync {
   std::mutex GlobalLock;
 };
 
+// const char* key hash for STL maps
+struct HashCStr {
+  size_t operator()(const char *S) const {
+    constexpr size_t Prime = 31;
+    size_t Res = 0;
+    char Ch = 0;
+
+    for (; (Ch = *S); S++) {
+      Res += Ch + (Prime * Res);
+    }
+    return Res;
+  }
+};
+
+// const char* key comparison for STL maps
+struct CmpCStr {
+  bool operator()(const char *A, const char *B) const {
+    return std::strcmp(A, B) == 0;
+  }
+};
+
+// Interface to iterate via C strings.
+class CStringIterator {
+public:
+  // Get the next string. Returns next string's pointer or nullptr.
+  virtual const char *next() = 0;
+};
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

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

@@ -0,0 +1,59 @@
+//==----- spec_constant.hpp - SYCL public experimental API header file -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// Based on:
+// https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md
+// TODO:
+// 1) implement the SPIRV interop part of the proposal
+// 2) support arbitrary spec constant type; only primitive types are
+//    supported currently
+// 3) move to the new upcoming spec constant specification (then 1 above is not
+//    needed)
+
+#pragma once
+
+#include <CL/sycl/detail/sycl_fe_intrins.hpp>
+#include <CL/sycl/exception.hpp>
+
+namespace cl {
+namespace sycl {
+namespace experimental {
+
+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__
+  spec_constant() {}
+#else
+  spec_constant(T Cst) : Val(Cst) {}
+#endif
+#ifndef __SYCL_DEVICE_ONLY__
+  T Val;
+#endif
+  friend class cl::sycl::program;
+
+public:
+  T get() const { // explicit access.
+#ifdef __SYCL_DEVICE_ONLY__
+    const char *TName = __unique_stable_name(ID);
+    return __sycl_getSpecConstantValue<T>(TName);
+#else
+    return Val;
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+  operator T() const { // implicit conversion.
+    return get();
+  }
+};
+
+} // namespace experimental
+} // namespace sycl
+} // namespace cl

sycl/include/CL/sycl/program.hpp

@@ -11,6 +11,7 @@
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
 #include <CL/sycl/detail/os_util.hpp>
+#include <CL/sycl/experimental/spec_constant.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/stl.hpp>
@@ -293,6 +294,25 @@ class program {
   /// \return the current state of this SYCL program.
   program_state get_state() const;
 
+  /// Set the value of the specialization constant identified by the 'ID' type
+  /// template parameter and return its instance.
+  /// \param cst the specialization constant value
+  /// \return a specialization constant instance corresponding to given type ID
+  ///         passed as a template parameter
+  template <typename ID, typename T>
+  experimental::spec_constant<T, ID> set_spec_constant(T Cst) {
+    constexpr const char *Name = detail::SpecConstantInfo<ID>::getName();
+    static_assert(std::is_integral<T>::value ||
+                      std::is_floating_point<T>::value,
+                  "unsupported specialization constant type");
+#ifdef __SYCL_DEVICE_ONLY__
+    return experimental::spec_constant<T, ID>();
+#else
+    set_spec_constant_impl(Name, &Cst, sizeof(T));
+    return experimental::spec_constant<T, ID>(Cst);
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
 private:
   program(shared_ptr_class<detail::program_impl> impl);
 
@@ -330,6 +350,8 @@ class program {
                               string_class buildOptions,
                               detail::OSModuleHandle M);
 
+  void set_spec_constant_impl(const char *Name, void *Data, size_t Size);
+
   shared_ptr_class<detail::program_impl> impl;
 
   template <class Obj>

sycl/source/CMakeLists.txt

@@ -103,6 +103,7 @@ set(SYCL_SOURCES
     "detail/scheduler/scheduler.cpp"
     "detail/scheduler/graph_processor.cpp"
     "detail/scheduler/graph_builder.cpp"
+    "detail/spec_constant_impl.cpp"
     "detail/sycl_mem_obj_t.cpp"
     "detail/usm/clusm.cpp"
     "detail/usm/usm_dispatch.cpp"

sycl/source/detail/os_util.cpp

@@ -57,6 +57,7 @@ struct ModuleInfo {
 };
 
 constexpr OSModuleHandle OSUtil::ExeModuleHandle;
+constexpr OSModuleHandle OSUtil::DummyModuleHandle;
 
 static int callback(struct dl_phdr_info *Info, size_t Size, void *Data) {
   auto Base = reinterpret_cast<unsigned char *>(Info->dlpi_addr);

sycl/source/detail/program_impl.cpp

@@ -8,6 +8,8 @@
 
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
+#include <CL/sycl/detail/pi.h>
+#include <CL/sycl/detail/spec_constant_impl.hpp>
 #include <CL/sycl/kernel.hpp>
 #include <detail/program_impl.hpp>
 
@@ -311,6 +313,7 @@ void program_impl::compile(const string_class &Options) {
   check_device_feature_support<info::device::is_compiler_available>(MDevices);
   vector_class<RT::PiDevice> Devices(get_pi_devices());
   const detail::plugin &Plugin = getPlugin();
+  ProgramManager::getInstance().flushSpecConstants(MProgram, *MContext);
   RT::PiResult Err = Plugin.call_nocheck<PiApiKind::piProgramCompile>(
       MProgram, Devices.size(), Devices.data(), Options.c_str(), 0, nullptr,
       nullptr, nullptr, nullptr);
@@ -329,6 +332,7 @@ void program_impl::build(const string_class &Options) {
   check_device_feature_support<info::device::is_compiler_available>(MDevices);
   vector_class<RT::PiDevice> Devices(get_pi_devices());
   const detail::plugin &Plugin = getPlugin();
+  ProgramManager::getInstance().flushSpecConstants(MProgram, *MContext);
   RT::PiResult Err = Plugin.call_nocheck<PiApiKind::piProgramBuild>(
       MProgram, Devices.size(), Devices.data(), Options.c_str(), nullptr,
       nullptr);
@@ -417,7 +421,8 @@ void program_impl::create_pi_program_with_kernel_name(
     OSModuleHandle Module, const string_class &KernelName) {
   assert(!MProgram && "This program already has an encapsulated PI program");
   ProgramManager &PM = ProgramManager::getInstance();
-  DeviceImage &Img = PM.getDeviceImage(Module, KernelName, get_context());
+  RTDeviceBinaryImage &Img =
+      PM.getDeviceImage(Module, KernelName, get_context());
   MProgram = PM.createPIProgram(Img, get_context());
 }
 
@@ -444,6 +449,13 @@ vector_class<device> program_impl::get_info<info::program::devices>() const {
   return get_devices();
 }
 
+void program_impl::set_spec_constant_impl(const char *Name, const void *ValAddr,
+                                          size_t ValSize) {
+  spec_constant_impl &SC =
+      ProgramManager::getInstance().resolveSpecConstant(this, Name);
+  SC.set(ValSize, ValAddr);
+}
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/program_impl.hpp

@@ -293,6 +293,16 @@ class program_impl {
   /// \return the current state of this SYCL program.
   program_state get_state() const { return MState; }
 
+  void set_spec_constant_impl(const char *Name, const void *ValAddr,
+                              size_t ValSize);
+
+  /// Returns the OS module handle this program belongs to. A program belongs to
+  /// an OS module if it was built from device image(s) belonging to that
+  /// module.
+  /// TODO Some programs can be linked from images belonging to different
+  ///      modules. May need a special fake handle for the resulting program.
+  OSModuleHandle getOSModuleHandle() const { return MProgramModuleHandle; }
+
 private:
   // Deligating Constructor used in Implementation.
   program_impl(ContextImplPtr Context, program_interop_handle_t InteropProgram,