[SYCL] Add SYCL 2020 info::device::built_in_kernel_ids support (#4996)

- Add `ProgramManager::getBuiltInKernelID`, which generates and caches built-in kernel IDs.
- Use that API to generate or look up built-in kernel IDs, when queried.
- Throw an exception in program manager when actually trying to use built-in kernels, since they are not yet fully supported.
- Add SYCL 2020 deprecation warning for `built_in_kernels` (old query).

sycl/include/CL/sycl/info/device_traits.def

@@ -58,6 +58,7 @@ __SYCL_PARAM_TRAITS_SPEC(device, is_linker_available, bool)
 __SYCL_PARAM_TRAITS_SPEC(device, execution_capabilities,
                   std::vector<info::execution_capability>)
 __SYCL_PARAM_TRAITS_SPEC(device, queue_profiling, bool)
+__SYCL_PARAM_TRAITS_SPEC(device, built_in_kernel_ids, std::vector<kernel_id>)
 __SYCL_PARAM_TRAITS_SPEC(device, built_in_kernels, std::vector<std::string>)
 __SYCL_PARAM_TRAITS_SPEC(device, platform, cl::sycl::platform)
 __SYCL_PARAM_TRAITS_SPEC(device, name, std::string)

sycl/include/CL/sycl/info/info_desc.hpp

@@ -20,6 +20,7 @@ class program;
 #endif
 class device;
 class platform;
+class kernel_id;
 
 // TODO: stop using OpenCL directly, use PI.
 namespace info {
@@ -109,7 +110,8 @@ enum class device : cl_device_info {
   is_linker_available = CL_DEVICE_LINKER_AVAILABLE,
   execution_capabilities = CL_DEVICE_EXECUTION_CAPABILITIES,
   queue_profiling = CL_DEVICE_QUEUE_PROPERTIES,
-  built_in_kernels = CL_DEVICE_BUILT_IN_KERNELS,
+  built_in_kernels __SYCL2020_DEPRECATED("use built_in_kernel_ids instead") =
+      CL_DEVICE_BUILT_IN_KERNELS,
   platform = CL_DEVICE_PLATFORM,
   name = CL_DEVICE_NAME,
   vendor = CL_DEVICE_VENDOR,
@@ -136,6 +138,7 @@ enum class device : cl_device_info {
   sub_group_sizes = CL_DEVICE_SUB_GROUP_SIZES_INTEL,
   partition_type_property,
   kernel_kernel_pipe_support,
+  built_in_kernel_ids,
   // USM
   usm_device_allocations = PI_USM_DEVICE_SUPPORT,
   usm_host_allocations = PI_USM_HOST_SUPPORT,

sycl/source/detail/device_impl.hpp

@@ -10,6 +10,7 @@
 
 #include <CL/sycl/aspects.hpp>
 #include <CL/sycl/detail/pi.hpp>
+#include <CL/sycl/kernel_bundle.hpp>
 #include <CL/sycl/stl.hpp>
 #include <detail/device_info.hpp>
 #include <detail/platform_impl.hpp>

sycl/source/detail/device_info.hpp

@@ -19,6 +19,7 @@
 #include <detail/platform_impl.hpp>
 #include <detail/platform_util.hpp>
 #include <detail/plugin.hpp>
+#include <detail/program_manager/program_manager.hpp>
 
 #include <chrono>
 #include <thread>
@@ -279,6 +280,25 @@ struct get_device_info<std::vector<info::execution_capability>,
   }
 };
 
+// Specialization for built in kernel identifiers
+template <>
+struct get_device_info<std::vector<kernel_id>,
+                       info::device::built_in_kernel_ids> {
+  static std::vector<kernel_id> get(RT::PiDevice dev, const plugin &Plugin) {
+    std::string result =
+        get_device_info<std::string, info::device::built_in_kernels>::get(
+            dev, Plugin);
+    auto names = split_string(result, ';');
+
+    std::vector<kernel_id> ids;
+    ids.reserve(names.size());
+    for (const auto &name : names) {
+      ids.push_back(ProgramManager::getInstance().getBuiltInKernelID(name));
+    }
+    return ids;
+  }
+};
+
 // Specialization for built in kernels, splits the string returned by OpenCL
 template <>
 struct get_device_info<std::vector<std::string>,
@@ -979,6 +999,12 @@ template <> inline bool get_device_info_host<info::device::queue_profiling>() {
   return true;
 }
 
+template <>
+inline std::vector<kernel_id>
+get_device_info_host<info::device::built_in_kernel_ids>() {
+  return {};
+}
+
 template <>
 inline std::vector<std::string>
 get_device_info_host<info::device::built_in_kernels>() {

sycl/source/detail/program_manager/program_manager.cpp

@@ -1341,6 +1341,19 @@ std::vector<kernel_id> ProgramManager::getAllSYCLKernelIDs() {
   return AllKernelIDs;
 }
 
+kernel_id ProgramManager::getBuiltInKernelID(const std::string &KernelName) {
+  std::lock_guard<std::mutex> BuiltInKernelIDsGuard(m_BuiltInKernelIDsMutex);
+
+  auto KernelID = m_BuiltInKernelIDs.find(KernelName);
+  if (KernelID == m_BuiltInKernelIDs.end()) {
+    auto Impl = std::make_shared<kernel_id_impl>(KernelName);
+    auto CachedID = createSyclObjFromImpl<kernel_id>(Impl);
+    KernelID = m_BuiltInKernelIDs.insert({KernelName, CachedID}).first;
+  }
+
+  return KernelID->second;
+}
+
 std::vector<device_image_plain>
 ProgramManager::getSYCLDeviceImagesWithCompatibleState(
     const context &Ctx, const std::vector<device> &Devs,
@@ -1511,6 +1524,17 @@ std::vector<device_image_plain> ProgramManager::getSYCLDeviceImages(
 std::vector<device_image_plain> ProgramManager::getSYCLDeviceImages(
     const context &Ctx, const std::vector<device> &Devs,
     const std::vector<kernel_id> &KernelIDs, bundle_state TargetState) {
+  {
+    std::lock_guard<std::mutex> BuiltInKernelIDsGuard(m_BuiltInKernelIDsMutex);
+
+    for (const kernel_id &ID : KernelIDs) {
+      if (m_BuiltInKernelIDs.find(ID.get_name()) != m_BuiltInKernelIDs.end())
+        throw sycl::exception(make_error_code(errc::kernel_argument),
+                              "Attempting to use a built-in kernel. They are "
+                              "not fully supported");
+    }
+  }
+
   // Collect device images with compatible state
   std::vector<device_image_plain> DeviceImages =
       getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, TargetState);

sycl/source/detail/program_manager/program_manager.hpp

@@ -178,6 +178,10 @@ class ProgramManager {
   // in SYCL device images.
   std::vector<kernel_id> getAllSYCLKernelIDs();
 
+  // The function returns the unique SYCL kernel identifier associated with a
+  // built-in kernel name.
+  kernel_id getBuiltInKernelID(const std::string &KernelName);
+
   // The function returns a vector of SYCL device images that are compiled with
   // the required state and at least one device from the passed list of devices.
   std::vector<device_image_plain>
@@ -327,6 +331,13 @@ class ProgramManager {
   /// Access must be guarded by the m_KernelIDsMutex mutex.
   std::unordered_set<std::string> m_ExportedSymbols;
 
+  /// Maps names of built-in kernels to their unique kernel IDs.
+  /// Access must be guarded by the m_BuiltInKernelIDsMutex mutex.
+  std::unordered_map<std::string, kernel_id> m_BuiltInKernelIDs;
+
+  /// Protects built-in kernel ID cache.
+  std::mutex m_BuiltInKernelIDsMutex;
+
   // Keeps track of pi_program to image correspondence. Needed for:
   // - knowing which specialization constants are used in the program and
   //   injecting their current values before compiling the SPIR-V; the binary

sycl/test/abi/sycl_symbols_linux.dump

@@ -4109,9 +4109,14 @@ _ZNK2cl4sycl6device3hasENS0_6aspectE
 _ZNK2cl4sycl6device6is_cpuEv
 _ZNK2cl4sycl6device6is_gpuEv
 _ZNK2cl4sycl6device7is_hostEv
+_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131072EEENS3_12param_traitsIS4_XT_EE11return_typeEv
+_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131073EEENS3_12param_traitsIS4_XT_EE11return_typeEv
+_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131074EEENS3_12param_traitsIS4_XT_EE11return_typeEv
+_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131075EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16648EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16649EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16650EEENS3_12param_traitsIS4_XT_EE11return_typeEv
+_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16651EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16784EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16785EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE16786EEENS3_12param_traitsIS4_XT_EE11return_typeEv
@@ -4244,6 +4249,7 @@ _ZNK2cl4sycl6streamneERKS1_
 _ZNK2cl4sycl7context11get_backendEv
 _ZNK2cl4sycl7context11get_devicesEv
 _ZNK2cl4sycl7context12get_platformEv
+_ZNK2cl4sycl7context12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
 _ZNK2cl4sycl7context12get_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEET_v
 _ZNK2cl4sycl7context12get_propertyINS0_8property5image12use_host_ptrEEET_v
 _ZNK2cl4sycl7context12get_propertyINS0_8property5image13context_boundEEET_v
@@ -4256,6 +4262,7 @@ _ZNK2cl4sycl7context12get_propertyINS0_8property6noinitEEET_v
 _ZNK2cl4sycl7context12get_propertyINS0_8property7context4cuda19use_primary_contextEEET_v
 _ZNK2cl4sycl7context12get_propertyINS0_8property7no_initEEET_v
 _ZNK2cl4sycl7context12get_propertyINS0_8property9reduction22initialize_to_identityEEET_v
+_ZNK2cl4sycl7context12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
 _ZNK2cl4sycl7context12has_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEEbv
 _ZNK2cl4sycl7context12has_propertyINS0_8property5image12use_host_ptrEEEbv
 _ZNK2cl4sycl7context12has_propertyINS0_8property5image13context_boundEEEbv
@@ -4268,8 +4275,6 @@ _ZNK2cl4sycl7context12has_propertyINS0_8property6noinitEEEbv
 _ZNK2cl4sycl7context12has_propertyINS0_8property7context4cuda19use_primary_contextEEEbv
 _ZNK2cl4sycl7context12has_propertyINS0_8property7no_initEEEbv
 _ZNK2cl4sycl7context12has_propertyINS0_8property9reduction22initialize_to_identityEEEbv
-_ZNK2cl4sycl7context12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
-_ZNK2cl4sycl7context12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
 _ZNK2cl4sycl7context3getEv
 _ZNK2cl4sycl7context7is_hostEv
 _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4224EEENS3_12param_traitsIS4_XT_EE11return_typeEv
@@ -4288,6 +4293,7 @@ _ZNK2cl4sycl7program11get_backendEv
 _ZNK2cl4sycl7program11get_contextEv
 _ZNK2cl4sycl7program11get_devicesEv
 _ZNK2cl4sycl7program12get_binariesEv
+_ZNK2cl4sycl7program12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property5image12use_host_ptrEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property5image13context_boundEEET_v
@@ -4298,9 +4304,9 @@ _ZNK2cl4sycl7program12get_propertyINS0_8property6buffer13context_boundEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property6buffer9use_mutexEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property6noinitEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property7context4cuda19use_primary_contextEEET_v
-_ZNK2cl4sycl7program12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property7no_initEEET_v
 _ZNK2cl4sycl7program12get_propertyINS0_8property9reduction22initialize_to_identityEEET_v
+_ZNK2cl4sycl7program12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property5image12use_host_ptrEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property5image13context_boundEEEbv
@@ -4311,7 +4317,6 @@ _ZNK2cl4sycl7program12has_propertyINS0_8property6buffer13context_boundEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property6buffer9use_mutexEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property6noinitEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property7context4cuda19use_primary_contextEEEbv
-_ZNK2cl4sycl7program12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property7no_initEEEbv
 _ZNK2cl4sycl7program12has_propertyINS0_8property9reduction22initialize_to_identityEEEbv
 _ZNK2cl4sycl7program16get_link_optionsB5cxx11Ev
@@ -4324,6 +4329,7 @@ _ZNK2cl4sycl7program8get_infoILNS0_4info7programE4449EEENS3_12param_traitsIS4_XT
 _ZNK2cl4sycl7program8get_infoILNS0_4info7programE4451EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl7program9getNativeEv
 _ZNK2cl4sycl7program9get_stateEv
+_ZNK2cl4sycl7sampler12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
 _ZNK2cl4sycl7sampler12get_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEET_v
 _ZNK2cl4sycl7sampler12get_propertyINS0_8property5image12use_host_ptrEEET_v
 _ZNK2cl4sycl7sampler12get_propertyINS0_8property5image13context_boundEEET_v
@@ -4336,7 +4342,7 @@ _ZNK2cl4sycl7sampler12get_propertyINS0_8property6noinitEEET_v
 _ZNK2cl4sycl7sampler12get_propertyINS0_8property7context4cuda19use_primary_contextEEET_v
 _ZNK2cl4sycl7sampler12get_propertyINS0_8property7no_initEEET_v
 _ZNK2cl4sycl7sampler12get_propertyINS0_8property9reduction22initialize_to_identityEEET_v
-_ZNK2cl4sycl7sampler12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v
+_ZNK2cl4sycl7sampler12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
 _ZNK2cl4sycl7sampler12has_propertyINS0_3ext6oneapi8property6buffer22use_pinned_host_memoryEEEbv
 _ZNK2cl4sycl7sampler12has_propertyINS0_8property5image12use_host_ptrEEEbv
 _ZNK2cl4sycl7sampler12has_propertyINS0_8property5image13context_boundEEEbv
@@ -4349,7 +4355,6 @@ _ZNK2cl4sycl7sampler12has_propertyINS0_8property6noinitEEEbv
 _ZNK2cl4sycl7sampler12has_propertyINS0_8property7context4cuda19use_primary_contextEEEbv
 _ZNK2cl4sycl7sampler12has_propertyINS0_8property7no_initEEEbv
 _ZNK2cl4sycl7sampler12has_propertyINS0_8property9reduction22initialize_to_identityEEEbv
-_ZNK2cl4sycl7sampler12has_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEEbv
 _ZNK2cl4sycl7sampler18get_filtering_modeEv
 _ZNK2cl4sycl7sampler19get_addressing_modeEv
 _ZNK2cl4sycl7sampler33get_coordinate_normalization_modeEv
@@ -4377,7 +4382,3 @@ _ZNK2cl4sycl9exception8categoryEv
 _ZNK2cl4sycl9kernel_id8get_nameEv
 __sycl_register_lib
 __sycl_unregister_lib
-_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131072EEENS3_12param_traitsIS4_XT_EE11return_typeEv
-_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131075EEENS3_12param_traitsIS4_XT_EE11return_typeEv
-_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131074EEENS3_12param_traitsIS4_XT_EE11return_typeEv
-_ZNK2cl4sycl6device8get_infoILNS0_4info6deviceE131073EEENS3_12param_traitsIS4_XT_EE11return_typeEv