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

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
+  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime")
   include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,7 @@
-# commit 0a90db9b2c36960c9b28ce18557ca15760724c4d
-# Merge: c4d9fdb4 6e0bdeb9
+# commit db83117e830406b0d9950e24892dba868acba354
+# Merge: 0a90db9b c79df596
 # Author: Callum Fare <[email protected]>
-# Date:   Wed Nov 27 12:16:44 2024 +0000
-#     Merge pull request #2369 from Bensuo/ben/kernel-binary-update-l0
-#     [CMDBUF] Implement kernel binary update for L0 adapter
-set(UNIFIED_RUNTIME_TAG 0a90db9b2c36960c9b28ce18557ca15760724c4d)
+# Date:   Wed Nov 27 16:04:19 2024 +0000
+#     Merge pull request #2261 from againull/againull/2d_block_exp
+#     Add new device descriptor to query 2D block array capabilities of the Intel GPU
+set(UNIFIED_RUNTIME_TAG db83117e830406b0d9950e24892dba868acba354)

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 a 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

@@ -1589,6 +1589,25 @@ 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) {
+  if (!Dev->has(aspect::ext_intel_esimd))
+    return false;
+
+  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) &&
+         (BlockArrayCapabilities &
+          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/InlineAsm/asm_glb.cpp

@@ -43,7 +43,8 @@ int main(void) {
     queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
     auto dev = q.get_device();
-    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+    std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+              << "\n";
 
     auto e = q.submit([&](handler &cgh) {
       auto PA = bufa.get_access<access::mode::read>(cgh);

sycl/test-e2e/ESIMD/addc.cpp

@@ -138,7 +138,7 @@ template <int N, bool AIsVector, bool BIsVector> bool test(sycl::queue Q) {
 int main() {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto D = Q.get_device();
-  std::cout << "Running on " << D.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << D.get_info<sycl::info::device::name>() << "\n";
 
   constexpr bool AIsVector = true;
   constexpr bool BIsVector = true;

sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp

@@ -46,7 +46,8 @@ int main(void) {
     queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
     auto dev = q.get_device();
-    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+    std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+              << "\n";
 
     auto e = q.submit([&](handler &cgh) {
       auto PA = bufa.get_access<access::mode::read>(cgh);

sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp

@@ -44,7 +44,8 @@ int main(void) {
     queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
     auto dev = q.get_device();
-    std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+    std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+              << "\n";
 
     auto e = q.submit([&](handler &cgh) {
       auto PA = bufa.get_access<access::mode::read>(cgh);

sycl/test-e2e/ESIMD/dpas/dpas_bf16.cpp

@@ -16,7 +16,8 @@
 int main(int argc, const char *argv[]) {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto Dev = Q.get_device();
-  std::cout << "Running on " << Dev.get_info<info::device::name>() << std::endl;
+  std::cout << "Running on " << Dev.get_info<sycl::info::device::name>()
+            << std::endl;
 
   bool Print = argc > 1 && std::string(argv[1]) == "-debug";
   bool Passed = true;

sycl/test-e2e/ESIMD/dpas/dpas_fp16.cpp

@@ -16,7 +16,8 @@
 int main(int argc, const char *argv[]) {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto Dev = Q.get_device();
-  std::cout << "Running on " << Dev.get_info<info::device::name>() << std::endl;
+  std::cout << "Running on " << Dev.get_info<sycl::info::device::name>()
+            << std::endl;
 
   bool Print = argc > 1 && std::string(argv[1]) == "-debug";
   bool Passed = true;

sycl/test-e2e/ESIMD/dpas/dpas_int.cpp

@@ -16,7 +16,8 @@
 int main(int argc, const char *argv[]) {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto Dev = Q.get_device();
-  std::cout << "Running on " << Dev.get_info<info::device::name>() << std::endl;
+  std::cout << "Running on " << Dev.get_info<sycl::info::device::name>()
+            << std::endl;
 
   bool Print = argc > 1 && std::string(argv[1]) == "-debug";
   bool Passed = true;

sycl/test-e2e/ESIMD/imulh_umulh.cpp

@@ -155,7 +155,7 @@ template <int N, bool AIsVector, bool BIsVector> bool tests(sycl::queue Q) {
 int main() {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto D = Q.get_device();
-  std::cout << "Running on " << D.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << D.get_info<sycl::info::device::name>() << "\n";
 
   constexpr bool AIsVector = true;
   constexpr bool BIsVector = true;

sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp

@@ -96,7 +96,8 @@ int main() {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   bool passed = true;
   passed &= test<char, 1>(q);

sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp

@@ -848,7 +848,8 @@ int main(void) {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   Config cfg{
       11,  // int threads_per_group;

sycl/test-e2e/ESIMD/lsc/local_accessor_atomic_smoke.cpp

@@ -627,7 +627,8 @@ int main(void) {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   bool passed = true;
 #ifndef CMPXCHG_TEST

sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp

@@ -102,7 +102,7 @@ int main() {
 
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto device = q.get_device();
-  std::cout << "Device name: " << device.get_info<info::device::name>()
+  std::cout << "Device name: " << device.get_info<sycl::info::device::name>()
             << std::endl;
 
   int error = testUSM<8>(q);

sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp

@@ -142,7 +142,7 @@ int main() {
   auto q =
       queue{esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()};
   auto device = q.get_device();
-  std::cout << "Device name: " << device.get_info<info::device::name>()
+  std::cout << "Device name: " << device.get_info<sycl::info::device::name>()
             << std::endl;
 
   int error = testUSM<8>(q);

sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp

@@ -89,7 +89,7 @@ int main() {
   auto q =
       queue{esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()};
   auto device = q.get_device();
-  std::cout << "Device name: " << device.get_info<info::device::name>()
+  std::cout << "Device name: " << device.get_info<sycl::info::device::name>()
             << std::endl;
 
   int error = testAccessor<8>(q);

sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_smoke.cpp

@@ -473,7 +473,8 @@ int main(void) {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   bool passed = true;
 #ifndef CMPXCHG_TEST

sycl/test-e2e/ESIMD/private_memory/private_memory.cpp

@@ -149,8 +149,8 @@ template <typename T> bool tests(queue Q) {
 
 int main() {
   queue Q;
-  std::cout << "Running on " << Q.get_device().get_info<info::device::name>()
-            << "\n";
+  std::cout << "Running on "
+            << Q.get_device().get_info<sycl::info::device::name>() << "\n";
 
   bool Passed = true;
   Passed &= tests<int8_t>(Q);

sycl/test-e2e/ESIMD/radix_sort.cpp

@@ -588,7 +588,8 @@ int main(int argc, char *argv[]) {
           property::queue::in_order());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
   auto ctxt = q.get_context();
 
   // allocate and initialized input

sycl/test-e2e/ESIMD/slm_alloc.cpp

@@ -92,7 +92,8 @@ __attribute__((noinline))
 int main(void) {
   queue q;
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
   std::cout << "force_inline=" << force_inline << "\n";
   auto ctxt = q.get_context();
   uint32_t size = SLM_TOTAL * NUM_WGS / ELEM_SIZE;

sycl/test-e2e/ESIMD/slm_alloc_many_kernels_many_funcs.cpp

@@ -111,7 +111,8 @@ INLINE_CTL void foo(int local_id, T *out, unsigned base) {
 int main(void) {
   queue q;
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
   std::cout << "force_inline=" << force_inline << "\n";
   auto ctxt = q.get_context();
 

sycl/test-e2e/ESIMD/slm_alloc_many_kernels_one_func.cpp

@@ -59,7 +59,8 @@ __attribute__((noinline))
 int main(void) {
   queue q;
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
   std::cout << "force_inline=" << force_inline << "\n";
   auto ctxt = q.get_context();
 

sycl/test-e2e/ESIMD/subb.cpp

@@ -140,7 +140,7 @@ template <int N, bool AIsVector, bool BIsVector> bool test(sycl::queue Q) {
 int main() {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto D = Q.get_device();
-  std::cout << "Running on " << D.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << D.get_info<sycl::info::device::name>() << "\n";
 
   constexpr bool AIsVector = true;
   constexpr bool BIsVector = true;

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

@@ -14,7 +14,8 @@ int main(void) {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   bool passed = true;
 

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

@@ -17,7 +17,8 @@ int main(void) {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   bool passed = true;
 

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

@@ -0,0 +1,39 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Test has_2d_block_io_supported device descriptor for some known
+// architectures.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/intel/esimd.hpp>
+
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue Q;
+  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-e2e/ESIMD/vadd_raw_send_gen12.cpp

@@ -35,7 +35,8 @@ int main(void) {
   queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
 
   auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
 
   int *A = malloc_shared<int>(Size, q);
   int *B = malloc_shared<int>(Size, q);

sycl/test-e2e/ESIMD/wait.cpp

@@ -57,7 +57,8 @@ bool test(sycl::queue Q, int IArg = 128) {
 int main() {
   queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
   auto Dev = Q.get_device();
-  std::cout << "Running on " << Dev.get_info<info::device::name>() << std::endl;
+  std::cout << "Running on " << Dev.get_info<sycl::info::device::name>()
+            << std::endl;
 
   bool Passed = true;
   Passed &= test(Q);