[SYCL] interim kernel compiler with SYCL support (#14172)

SYCL language support on the part of the kernel_compiler is specified
here: #11985

However, that specification is not presently realizable. We need more
support from the FE and post link tool to get the demangled names
propagated through. But it is usable before that, with constraints about
using extern "C" or knowing the mangled kernel name. We have folk that
want to test in the interim. I've refrained from updating the spec, and
instead this interim support in our experimental extension will be snuck
in until it can be completed in full - then we'll update the spec to
release it properly.

Author: cperkinsintel

clang/lib/Sema/SemaSYCL.cpp

@@ -1111,18 +1111,24 @@ static std::pair<std::string, std::string> constructFreeFunctionKernelName(
     SemaSYCL &SemaSYCLRef, const FunctionDecl *FreeFunc, MangleContext &MC) {
   SmallString<256> Result;
   llvm::raw_svector_ostream Out(Result);
+  std::string NewName;
   std::string StableName;
 
-  MC.mangleName(FreeFunc, Out);
-  std::string MangledName(Out.str());
-  size_t StartNums = MangledName.find_first_of("0123456789");
-  size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums);
-  size_t NameLength =
-      std::stoi(MangledName.substr(StartNums, EndNums - StartNums));
-  size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength;
-  std::string NewName = MangledName.substr(0, StartNums) +
-                        std::to_string(NewNameLength) + "__sycl_kernel_" +
-                        MangledName.substr(EndNums);
+  // Handle extern "C"
+  if (FreeFunc->getLanguageLinkage() == CLanguageLinkage) {
+    const IdentifierInfo *II = FreeFunc->getIdentifier();
+    NewName = "__sycl_kernel_" + II->getName().str();
+  } else {
+    MC.mangleName(FreeFunc, Out);
+    std::string MangledName(Out.str());
+    size_t StartNums = MangledName.find_first_of("0123456789");
+    size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums);
+    size_t NameLength =
+        std::stoi(MangledName.substr(StartNums, EndNums - StartNums));
+    size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength;
+    NewName = MangledName.substr(0, StartNums) + std::to_string(NewNameLength) +
+              "__sycl_kernel_" + MangledName.substr(EndNums);
+  }
   StableName = NewName;
   return {NewName, StableName};
 }

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -209,8 +209,10 @@ enum PropKind : uint32_t {
   CallsIndirectly = 68,
   InputDataPlacement = 69,
   OutputDataPlacement = 70,
+  IncludeFiles = 71,
+  RegisteredKernelNames = 72,
   // PropKindSize must always be the last value.
-  PropKindSize = 71,
+  PropKindSize = 73,
 };
 
 struct property_key_base_tag {};

sycl/include/sycl/kernel_bundle.hpp

@@ -814,6 +814,32 @@ build(const kernel_bundle<bundle_state::input> &InputBundle,
 
 namespace ext::oneapi::experimental {
 
+namespace detail {
+struct create_bundle_from_source_props;
+struct build_source_bundle_props;
+} // namespace detail
+
+/////////////////////////
+// PropertyT syclex::include_files
+/////////////////////////
+struct include_files
+    : detail::run_time_property_key<detail::PropKind::IncludeFiles> {
+  include_files();
+  include_files(const std::string &name, const std::string &content) {
+    record.emplace_back(std::make_pair(name, content));
+  }
+  void add(const std::string &name, const std::string &content) {
+    record.emplace_back(std::make_pair(name, content));
+  }
+  std::vector<std::pair<std::string, std::string>> record;
+};
+using include_files_key = include_files;
+
+template <>
+struct is_property_key_of<include_files_key,
+                          detail::create_bundle_from_source_props>
+    : std::true_type {};
+
 /////////////////////////
 // PropertyT syclex::build_options
 /////////////////////////
@@ -826,8 +852,7 @@ struct build_options
 using build_options_key = build_options;
 
 template <>
-struct is_property_key_of<build_options_key,
-                          sycl::kernel_bundle<bundle_state::ext_oneapi_source>>
+struct is_property_key_of<build_options_key, detail::build_source_bundle_props>
     : std::true_type {};
 
 /////////////////////////
@@ -840,72 +865,132 @@ struct save_log : detail::run_time_property_key<detail::PropKind::BuildLog> {
 using save_log_key = save_log;
 
 template <>
-struct is_property_key_of<save_log_key,
-                          sycl::kernel_bundle<bundle_state::ext_oneapi_source>>
+struct is_property_key_of<save_log_key, detail::build_source_bundle_props>
     : std::true_type {};
 
 /////////////////////////
-// syclex::is_source_kernel_bundle_supported
+// PropertyT syclex::registered_kernel_names
 /////////////////////////
+struct registered_kernel_names
+    : detail::run_time_property_key<detail::PropKind::RegisteredKernelNames> {
+  std::vector<std::string> kernel_names;
+  registered_kernel_names() {}
+  registered_kernel_names(const std::string &knArg) : kernel_names{knArg} {}
+  registered_kernel_names(const std::vector<std::string> &knsArg)
+      : kernel_names(knsArg) {}
+  void add(const std::string &name) { kernel_names.push_back(name); }
+};
+using registered_kernel_names_key = registered_kernel_names;
+
+template <>
+struct is_property_key_of<registered_kernel_names_key,
+                          detail::build_source_bundle_props> : std::true_type {
+};
+
+namespace detail {
+// forward decls
 __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE,
                                                      source_language Language);
 
-/////////////////////////
-// syclex::create_kernel_bundle_from_source
-/////////////////////////
-
 __SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
-create_kernel_bundle_from_source(const context &SyclContext,
-                                 source_language Language,
-                                 const std::string &Source);
+make_kernel_bundle_from_source(
+    const context &SyclContext, source_language Language,
+    const std::string &Source,
+    std::vector<std::pair<std::string, std::string>> IncludePairsVec);
 
 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
 __SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
-create_kernel_bundle_from_source(const context &SyclContext,
-                                 source_language Language,
-                                 const std::vector<std::byte> &Bytes);
+make_kernel_bundle_from_source(
+    const context &SyclContext, source_language Language,
+    const std::vector<std::byte> &Bytes,
+    std::vector<std::pair<std::string, std::string>> IncludePairsVec);
 #endif
 
-/////////////////////////
-// syclex::build(source_kb) => exe_kb
-/////////////////////////
-namespace detail {
-// forward decl
 __SYCL_EXPORT kernel_bundle<bundle_state::executable>
 build_from_source(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
                   const std::vector<device> &Devices,
                   const std::vector<std::string> &BuildOptions,
-                  std::string *LogPtr);
+                  std::string *LogPtr,
+                  const std::vector<std::string> &RegisteredKernelNames);
 
 } // namespace detail
 
+/////////////////////////
+// syclex::create_kernel_bundle_from_source
+/////////////////////////
+template <
+    typename PropertyListT = empty_properties_t,
+    typename = std::enable_if_t<
+        is_property_list_v<PropertyListT> &&
+        detail::all_props_are_keys_of<detail::create_bundle_from_source_props,
+                                      PropertyListT>::value>>
+kernel_bundle<bundle_state::ext_oneapi_source> create_kernel_bundle_from_source(
+    const context &SyclContext, source_language Language,
+    const std::string &Source, PropertyListT props = {}) {
+  std::vector<std::pair<std::string, std::string>> IncludePairsVec;
+  if constexpr (props.template has_property<include_files>()) {
+    IncludePairsVec = props.template get_property<include_files>().record;
+  }
+
+  return detail::make_kernel_bundle_from_source(SyclContext, Language, Source,
+                                                IncludePairsVec);
+}
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <
+    typename PropertyListT = empty_properties_t,
+    typename = std::enable_if_t<
+        is_property_list_v<PropertyListT> &&
+        detail::all_props_are_keys_of<detail::create_bundle_from_source_props,
+                                      PropertyListT>::value>>
+kernel_bundle<bundle_state::ext_oneapi_source> create_kernel_bundle_from_source(
+    const context &SyclContext, source_language Language,
+    const std::vector<std::byte> &Bytes, PropertyListT props = {}) {
+  std::vector<std::pair<std::string, std::string>> IncludePairsVec;
+  if constexpr (props.template has_property<include_files>()) {
+    IncludePairsVec = props.template get_property<include_files>().record;
+  }
+
+  return detail::make_kernel_bundle_from_source(SyclContext, Language, Bytes,
+                                                IncludePairsVec);
+}
+#endif
+
+/////////////////////////
+// syclex::build(source_kb) => exe_kb
+/////////////////////////
+
 template <typename PropertyListT = empty_properties_t,
           typename = std::enable_if_t<
               is_property_list_v<PropertyListT> &&
-              detail::all_props_are_keys_of<
-                  kernel_bundle<bundle_state::ext_oneapi_source>,
-                  PropertyListT>::value>>
+              detail::all_props_are_keys_of<detail::build_source_bundle_props,
+                                            PropertyListT>::value>>
 
 kernel_bundle<bundle_state::executable>
 build(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
       const std::vector<device> &Devices, PropertyListT props = {}) {
   std::vector<std::string> BuildOptionsVec;
   std::string *LogPtr = nullptr;
+  std::vector<std::string> RegisteredKernelNamesVec;
   if constexpr (props.template has_property<build_options>()) {
     BuildOptionsVec = props.template get_property<build_options>().opts;
   }
   if constexpr (props.template has_property<save_log>()) {
     LogPtr = props.template get_property<save_log>().log;
   }
-  return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr);
+  if constexpr (props.template has_property<registered_kernel_names>()) {
+    RegisteredKernelNamesVec =
+        props.template get_property<registered_kernel_names>().kernel_names;
+  }
+  return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr,
+                                   RegisteredKernelNamesVec);
 }
 
 template <typename PropertyListT = empty_properties_t,
           typename = std::enable_if_t<
               is_property_list_v<PropertyListT> &&
-              detail::all_props_are_keys_of<
-                  kernel_bundle<bundle_state::ext_oneapi_source>,
-                  PropertyListT>::value>>
+              detail::all_props_are_keys_of<detail::build_source_bundle_props,
+                                            PropertyListT>::value>>
 kernel_bundle<bundle_state::executable>
 build(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
       PropertyListT props = {}) {

sycl/include/sycl/kernel_bundle_enums.hpp

@@ -20,7 +20,7 @@ enum class bundle_state : char {
 
 namespace ext::oneapi::experimental {
 
-enum class source_language : int { opencl = 0, spirv = 1 /* sycl, cuda */ };
+enum class source_language : int { opencl = 0, spirv = 1, sycl = 2 /* cuda */ };
 
 // opencl versions
 struct cl_version {

sycl/source/CMakeLists.txt

@@ -217,6 +217,7 @@ set(SYCL_COMMON_SOURCES
     "detail/jit_compiler.cpp"
     "detail/jit_device_binaries.cpp"
     "detail/kernel_compiler/kernel_compiler_opencl.cpp"
+    "detail/kernel_compiler/kernel_compiler_sycl.cpp"
     "detail/kernel_impl.cpp"
     "detail/kernel_program_cache.cpp"
     "detail/memory_manager.cpp"

sycl/source/detail/device_impl.cpp

@@ -824,7 +824,8 @@ bool device_impl::isGetDeviceAndHostTimerSupported() {
 bool device_impl::extOneapiCanCompile(
     ext::oneapi::experimental::source_language Language) {
   try {
-    return is_source_kernel_bundle_supported(getBackend(), Language);
+    return sycl::ext::oneapi::experimental::detail::
+        is_source_kernel_bundle_supported(getBackend(), Language);
   } catch (sycl::exception &) {
     return false;
   }