[SYCL] interim kernel compiler with SYCL support

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

@@ -862,7 +862,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;
   }