[SYCL][Fusion] Add fusion device info descriptor (#8254)

Implement the
`sycl::ext::codeplay::experimental::info::device::supports_fusion`
device descriptor defined by the [kernel fusion extension
proposal](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc).

Kernel fusion is currently limited to SPIR-V based backends, so support
can be determined based on the device's backend.

---------

Signed-off-by: Lukas Sommer <[email protected]>

Author: sommerlukas

sycl/include/sycl/detail/info_desc_helpers.hpp

@@ -113,6 +113,7 @@ struct IsSubGroupInfo<info::kernel_device_specific::compile_sub_group_size>
       : std::true_type {                                                       \
     using return_type = Namespace::info::DescType::Desc::return_type;          \
   };
+#include <sycl/info/ext_codeplay_device_traits.def>
 #include <sycl/info/ext_intel_device_traits.def>
 #include <sycl/info/ext_oneapi_device_traits.def>
 #undef __SYCL_PARAM_TRAITS_SPEC

sycl/include/sycl/detail/pi.h

@@ -322,6 +322,7 @@ typedef enum {
   PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D = 0x20002,
   PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D = 0x20003,
   PI_EXT_ONEAPI_DEVICE_INFO_CUDA_ASYNC_BARRIER = 0x20004,
+  PI_EXT_CODEPLAY_DEVICE_INFO_SUPPORTS_FUSION = 0x20005,
 } _pi_device_info;
 
 typedef enum {

sycl/include/sycl/info/ext_codeplay_device_traits.def

@@ -0,0 +1,9 @@
+#ifndef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
+#define __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
+#define __SYCL_PARAM_TRAITS_TEMPLATE_SPEC __SYCL_PARAM_TRAITS_SPEC
+#endif
+__SYCL_PARAM_TRAITS_SPEC(ext::codeplay::experimental,device, supports_fusion, bool, PI_EXT_CODEPLAY_DEVICE_INFO_SUPPORTS_FUSION)
+#ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
+#undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
+#undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
+#endif

sycl/include/sycl/info/info_desc.hpp

@@ -184,6 +184,7 @@ template <typename T, T param> struct compatibility_param_traits {};
 namespace ext::oneapi::experimental::info::device {
 template <int Dimensions> struct max_work_groups;
 } // namespace ext::oneapi::experimental::info::device
+#include <sycl/info/ext_codeplay_device_traits.def>
 #include <sycl/info/ext_intel_device_traits.def>
 #include <sycl/info/ext_oneapi_device_traits.def>
 #undef __SYCL_PARAM_TRAITS_SPEC

sycl/source/detail/device_info.hpp

@@ -17,6 +17,7 @@
 #include <sycl/detail/os_util.hpp>
 #include <sycl/detail/pi.hpp>
 #include <sycl/device.hpp>
+#include <sycl/feature_test.hpp>
 #include <sycl/info/info_desc.hpp>
 #include <sycl/memory_enums.hpp>
 #include <sycl/platform.hpp>
@@ -751,6 +752,25 @@ struct get_device_info_impl<bool, info::device::ext_intel_mem_channel> {
   }
 };
 
+// Specialization for kernel fusion support
+template <>
+struct get_device_info_impl<
+    bool, ext::codeplay::experimental::info::device::supports_fusion> {
+  static bool get(RT::PiDevice dev, const plugin &Plugin) {
+#if SYCL_EXT_CODEPLAY_KERNEL_FUSION
+    // Currently fusion is only supported for SPIR-V based backends, i.e. OpenCL
+    // and LevelZero.
+    (void)dev;
+    return (Plugin.getBackend() == backend::ext_oneapi_level_zero) ||
+           (Plugin.getBackend() == backend::opencl);
+#else  // SYCL_EXT_CODEPLAY_KERNEL_FUSION
+    (void)dev;
+    (void)Plugin;
+    return false;
+#endif // SYCL_EXT_CODEPLAY_KERNEL_FUSION
+  }
+};
+
 template <typename Param>
 typename Param::return_type get_device_info(RT::PiDevice dev,
                                             const plugin &Plugin) {
@@ -1609,6 +1629,13 @@ get_device_info_host<ext::intel::info::device::max_compute_queue_indices>() {
       PI_ERROR_INVALID_DEVICE);
 }
 
+template <>
+inline bool get_device_info_host<
+    ext::codeplay::experimental::info::device::supports_fusion>() {
+  // No support for fusion on the host device.
+  return false;
+}
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/source/device.cpp

@@ -190,6 +190,7 @@ device::get_info<info::device::aspects>() const {
   template __SYCL_EXPORT ReturnT                                               \
   device::get_info<Namespace::info::DescType::Desc>() const;
 
+#include <sycl/info/ext_codeplay_device_traits.def>
 #include <sycl/info/ext_intel_device_traits.def>
 #include <sycl/info/ext_oneapi_device_traits.def>
 #undef __SYCL_PARAM_TRAITS_SPEC

sycl/test/abi/sycl_symbols_linux.dump

@@ -4242,6 +4242,7 @@ _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_wor
 _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi3EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device22max_global_work_groupsEEENS0_6detail19is_device_info_descIT_E11return_typeEv
+_ZNK4sycl3_V16device8get_infoINS0_3ext8codeplay12experimental4info6device15supports_fusionEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device10extensionsEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device11device_typeEEENS0_6detail19is_device_info_descIT_E11return_typeEv
 _ZNK4sycl3_V16device8get_infoINS0_4info6device12address_bitsEEENS0_6detail19is_device_info_descIT_E11return_typeEv

sycl/test/abi/sycl_symbols_windows.dump

@@ -162,6 +162,7 @@
 ??$get_info@Usingle_fp_config@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA?AV?$vector@W4fp_config@info@_V1@sycl@@V?$allocator@W4fp_config@info@_V1@sycl@@@std@@@std@@XZ
 ??$get_info@Usub_group_independent_forward_progress@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ
 ??$get_info@Usub_group_sizes@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA?AV?$vector@_KV?$allocator@_K@std@@@std@@XZ
+??$get_info@Usupports_fusion@device@info@experimental@codeplay@ext@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ
 ??$get_info@Uusm_device_allocations@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ
 ??$get_info@Uusm_host_allocations@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ
 ??$get_info@Uusm_restricted_shared_allocations@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ