[SYCL] Add esimd device descriptor for 2d load/store

Author: againull

sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md

@@ -31,6 +31,7 @@
   - [__regcall Calling convention](#__regcall-calling-convention)
   - [Inline assembly](#inline-assembly)
   - [Device aspect](#device-aspect)
+  - [Device Information Descriptors](#device-information-descriptors)
 - [Device queries and conditional dispatching of the code](#device-queries-and-conditional-dispatching-of-the-code)
 - [Implementation restrictions](#implementation-restrictions)
   - [Features not supported with the ESIMD extension](#features-not-supported-with-the-esimd-extension)
@@ -1018,6 +1019,11 @@ The new aspect has the following behavior when queried via `device::has()`:
 |--------|-------------|
 |`aspect::ext_intel_esimd` | Indicates that the device supports the `sycl_ext_intel_esimd` extension as defined in this document. |
 
+## Device Information Descriptors
+| Device Descriptors | Return Type | Description |
+| ------------------ | ----------- | ----------- |
+| `ext::intel::esimd::info::device::has_2d_block_io_support` | bool | Returns the boolean indicating whether 2D load/store/prefetch instructions are supported by the device. |
+
 ## Examples
 ### Vector addition (USM)
 ```cpp

sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md

@@ -552,7 +552,7 @@ Loads and returns a vector `simd<T, N>` where `N` is `BlockWidth * BlockHeight *
 `props` - The optional compile-time properties. Only cache hint properties are used.
 
 ### Restrictions
-* This function is available only for Intel® Data Center GPU Max Series (aka PVC).
+* This function is available only for devices with the `ext::intel::esimd::info::device::has_2d_block_io_support` information descriptor value equal to `true`.
 * `Cache-hint` properties, if passed, must follow the [rules](#valid-combinations-of-l1-and-l2-cache-hints-for-load-functions) for `load` functions.
 * `Transformed` and `Transposed` cannot be set to true at the same time.
 * `BlockWidth` * `BlockHeight` * `NBlocks` * sizeof(`T`) must not exceed 2048.
@@ -598,7 +598,7 @@ Prefetches elements from a memory block of the size `BlockWidth * BlockHeight *
 `props` - The compile-time properties, which must specify cache-hints.
 
 ### Restrictions
-* This function is available only for Intel® Data Center GPU Max Series (aka PVC).
+* This function is available only for devices with the `ext::intel::esimd::info::device::has_2d_block_io_support` information descriptor value equal to `true`.
 * `Cache-hint` properties must follow the [rules](#valid-combinations-of-l1-and-l2-cache-hints-for-prefetch-functions) for `prefetch` functions.
 * `BlockWidth` * `BlockHeight` * `NBlocks` * sizeof(`T`) must not exceed 2048.
 * `NBlocks` must be {1,2,4} for `bytes` and `words`, {1,2} for `dwords`, 1 for `qwords`.
@@ -630,7 +630,7 @@ Stores the vector `Vals` of the type `simd<T, N>` to 2D memory block where `N` i
 `props` - The optional compile-time properties. Only cache hint properties are used.
 
 ### Restrictions
-* This function is available only for Intel® Data Center GPU Max Series (aka PVC).
+* This function is available only for devices with the `ext::intel::esimd::info::device::has_2d_block_io_support` information descriptor value equal to `true`.
 * `Cache-hint` properties, if passed, must follow the [rules](#valid-combinations-of-l1-and-l2-cache-hints-for-store-functions) for `store` functions.
 * `BlockWidth` * `BlockHeight` * sizeof(`T`) must not exceed 512.
 * `BlockHeight` must not exceed 8.

sycl/include/sycl/info/ext_intel_device_traits.def

@@ -16,6 +16,7 @@ __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, free_memory, uint64_t, UR_DEVICE_IN
 __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, memory_clock_rate, uint32_t, UR_DEVICE_INFO_MEMORY_CLOCK_RATE)
 __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, memory_bus_width, uint32_t, UR_DEVICE_INFO_MEMORY_BUS_WIDTH)
 __SYCL_PARAM_TRAITS_SPEC(ext::intel, device, max_compute_queue_indices, int32_t, UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES)
+__SYCL_PARAM_TRAITS_SPEC(ext::intel::esimd, device, has_2d_block_io_support, bool, UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP)
 #ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF

sycl/source/detail/device_info.hpp

@@ -1563,6 +1563,21 @@ get_device_info<ext::intel::info::device::memory_bus_width>(
   return get_device_info_impl<Param::return_type, Param>::get(Dev);
 }
 
+template <>
+inline ext::intel::esimd::info::device::has_2d_block_io_support::return_type
+get_device_info<ext::intel::esimd::info::device::has_2d_block_io_support>(
+    const DeviceImplPtr &Dev) {
+  ur_exp_device_2d_block_array_capability_flags_t BlockArrayCapabilities;
+  Dev->getAdapter()->call<UrApiKind::urDeviceGetInfo>(
+      Dev->getHandleRef(),
+      UrInfoCode<
+          ext::intel::esimd::info::device::has_2d_block_io_support>::value,
+      sizeof(BlockArrayCapabilities), &BlockArrayCapabilities, nullptr);
+  return BlockArrayCapabilities &
+         (UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD |
+          UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE);
+}
+
 // Returns the list of all progress guarantees that can be requested for
 // work_groups from the coordination level of root_group when using the device
 // given by Dev. First it calls getProgressGuarantee to get the strongest

sycl/test-e2e/ESIMD/unified_memory_api/device_queries.cpp

@@ -0,0 +1,41 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Test to check has_2d_block_io_supported device descriptor.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/intel/esimd.hpp>
+
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue Q(sycl::gpu_selector_v);
+
+  // if architecture is pvc, then has_2d_block_io_support is expected to be
+  // true.
+  auto Arch = Q.get_device().get_info<syclex::info::device::architecture>();
+  bool Has2DBlockIOSupport =
+      Q.get_device()
+          .get_info<
+              sycl::ext::intel::esimd::info::device::has_2d_block_io_support>();
+  if (Arch == syclex::architecture::intel_gpu_pvc) {
+    if (!Has2DBlockIOSupport) {
+      std::cerr << "Error: has_2d_block_io_support is expected to be true for "
+                   "PVC architecture"
+                << std::endl;
+      return 1;
+    }
+  }
+  if (Arch == syclex::architecture::intel_gpu_tgllp ||
+      Arch == syclex::architecture::intel_gpu_dg2_g10 ||
+      Arch == syclex::architecture::intel_gpu_dg2_g11 ||
+      Arch == syclex::architecture::intel_gpu_dg2_g12) {
+    if (Has2DBlockIOSupport) {
+      std::cerr << "Error: has_2d_block_io_support is expected to be false for "
+                   "Tiger Lake and DG2"
+                << std::endl;
+      return 1;
+    }
+  }
+  return 0;
+}

sycl/test/abi/sycl_symbols_linux.dump

@@ -3242,6 +3242,10 @@ _ZN4sycl3_V16detail13lgamma_r_implEfPi
 _ZN4sycl3_V16detail13make_platformEmNS0_7backendE
 _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE
 _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE
+_ZN4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv
+_ZN4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv
+_ZN4sycl3_V16detail14SubmissionInfoC1Ev
+_ZN4sycl3_V16detail14SubmissionInfoC2Ev
 _ZN4sycl3_V16detail14addCounterInitERNS0_7handlerERSt10shared_ptrINS1_10queue_implEERS4_IiE
 _ZN4sycl3_V16detail14getBorderColorENS0_19image_channel_orderE
 _ZN4sycl3_V16detail14tls_code_loc_t5queryEv
@@ -3251,10 +3255,6 @@ _ZN4sycl3_V16detail14tls_code_loc_tC2ERKNS1_13code_locationE
 _ZN4sycl3_V16detail14tls_code_loc_tC2Ev
 _ZN4sycl3_V16detail14tls_code_loc_tD1Ev
 _ZN4sycl3_V16detail14tls_code_loc_tD2Ev
-_ZN4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv
-_ZN4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv
-_ZN4sycl3_V16detail14SubmissionInfoC1Ev
-_ZN4sycl3_V16detail14SubmissionInfoC2Ev
 _ZN4sycl3_V16detail16AccessorBaseHost10getAccDataEv
 _ZN4sycl3_V16detail16AccessorBaseHost14getAccessRangeEv
 _ZN4sycl3_V16detail16AccessorBaseHost14getMemoryRangeEv
@@ -3675,6 +3675,7 @@ _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext5intel4info6device25gpu_eu_co
 _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext5intel4info6device25max_compute_queue_indicesEEENT_11return_typeEv
 _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext5intel4info6device4uuidEEENT_11return_typeEv
 _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext5intel4info6device9device_idEEENT_11return_typeEv
+_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext5intel5esimd4info6device23has_2d_block_io_supportEEENT_11return_typeEv
 _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device12architectureEEENT_11return_typeEv
 _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi1EEEEENT_11return_typeEv
 _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENT_11return_typeEv
@@ -3780,6 +3781,7 @@ _ZNK4sycl3_V16device13get_info_implINS0_3ext5intel4info6device25gpu_eu_count_per
 _ZNK4sycl3_V16device13get_info_implINS0_3ext5intel4info6device25max_compute_queue_indicesEEENS0_6detail11ABINeutralTINS8_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext5intel4info6device4uuidEEENS0_6detail11ABINeutralTINS8_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext5intel4info6device9device_idEEENS0_6detail11ABINeutralTINS8_19is_device_info_descIT_E11return_typeEE4typeEv
+_ZNK4sycl3_V16device13get_info_implINS0_3ext5intel5esimd4info6device23has_2d_block_io_supportEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device12architectureEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi1EEEEENS0_6detail11ABINeutralTINSA_19is_device_info_descIT_E11return_typeEE4typeEv
 _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENS0_6detail11ABINeutralTINSA_19is_device_info_descIT_E11return_typeEE4typeEv

sycl/unittests/kernel-and-program/DeviceInfo.cpp

@@ -72,6 +72,16 @@ static ur_result_t redefinedDeviceGetInfo(void *pParams) {
     }
   }
 
+  if (*params.ppropName == UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP) {
+    assert(*params.ppropSize ==
+           sizeof(ur_exp_device_2d_block_array_capability_flags_t));
+    if (*params.ppPropValue) {
+      *static_cast<ur_exp_device_2d_block_array_capability_flags_t *>(
+          *params.ppPropValue) =
+          UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD |
+          UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE;
+    }
+  }
   return UR_RESULT_SUCCESS;
 }
 
@@ -185,6 +195,17 @@ TEST_F(DeviceInfoTest, GetDeviceMemoryBusWidth) {
       << "Expect memory_bus_width to be of uint32_t size";
 }
 
+TEST_F(DeviceInfoTest, GetDeviceESIMD2DBlockIOSupport) {
+  context Ctx{Plt.get_devices()[0]};
+  TestContext.reset(new TestCtx(Ctx));
+
+  device Dev = Ctx.get_devices()[0];
+
+  auto HasSupport =
+      Dev.get_info<ext::intel::esimd::info::device::has_2d_block_io_support>();
+  EXPECT_TRUE(HasSupport);
+}
+
 TEST_F(DeviceInfoTest, BuiltInKernelIDs) {
   context Ctx{Plt.get_devices()[0]};
   TestContext.reset(new TestCtx(Ctx));