Make this a single esimd aspect for 2d opeations

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -86,8 +86,7 @@ def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group"
 def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
 def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">;
 def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">;
-def AspectExt_intel_2d_block_load : Aspect<"ext_intel_2d_block_load">;
-def AspectExt_intel_2d_block_store : Aspect<"ext_intel_2d_block_store">;
+def AspectExt_intel_esimd_2d_block_io : Aspect<"ext_intel_esimd_2d_block_io">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -153,8 +152,7 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_intel_fpga_task_sequence,
     AspectExt_oneapi_atomic16,
     AspectExt_oneapi_virtual_functions,
-    AspectExt_intel_2d_block_load,
-    AspectExt_intel_2d_block_store],
+    AspectExt_intel_esimd_2d_block_io],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

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/againull/unified-runtime")
+  set(UNIFIED_RUNTIME_REPO "/iusers/againull/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 +1 @@
-set(UNIFIED_RUNTIME_TAG e5e7feccd5bb5fdae79cbbe51e60b4fe61e50598)
+set(UNIFIED_RUNTIME_TAG 17ae9823f997c4fddcb6ef5d1bbb0b843ec29220)

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 on devices with `ext_intel_esimd_2d_block_io` aspect.
 * `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 on devices with `ext_intel_esimd_2d_block_io` aspect.
 * `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 on devices with `ext_intel_esimd_2d_block_io` aspect.
 * `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/ext/intel/esimd/memory.hpp

@@ -13844,6 +13844,7 @@ template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1,
           int N = detail::get_lsc_block_2d_data_size<
               T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(),
           typename PropertyListT = oneapi::experimental::empty_properties_t>
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_intel_esimd_2d_block_io)]]
 __ESIMD_API std::enable_if_t<
     ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
 load_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,
@@ -13887,6 +13888,7 @@ template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1,
               T, NBlocks, BlockHeight, BlockWidth, false /*Transposed*/,
               false /*Transformed*/>(),
           typename PropertyListT = oneapi::experimental::empty_properties_t>
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_intel_esimd_2d_block_io)]]
 __ESIMD_API std::enable_if_t<
     ext::oneapi::experimental::is_property_list_v<PropertyListT>>
 prefetch_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,
@@ -13923,6 +13925,7 @@ template <typename T, int BlockWidth, int BlockHeight = 1,
               T, 1u, BlockHeight, BlockWidth, false /*Transposed*/,
               false /*Transformed*/>(),
           typename PropertyListT = oneapi::experimental::empty_properties_t>
+[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_intel_esimd_2d_block_io)]]
 __ESIMD_API std::enable_if_t<
     ext::oneapi::experimental::is_property_list_v<PropertyListT>>
 store_2d(T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,

sycl/include/sycl/info/aspects.def

@@ -72,5 +72,4 @@ __SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 78)
 __SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79)
 __SYCL_ASPECT(ext_oneapi_atomic16, 80)
 __SYCL_ASPECT(ext_oneapi_virtual_functions, 81)
-__SYCL_ASPECT(ext_intel_2d_block_load, 82)
-__SYCL_ASPECT(ext_intel_2d_block_store, 83)
+__SYCL_ASPECT(ext_intel_esimd_2d_block_io, 82)

sycl/source/detail/device_impl.cpp

@@ -777,7 +777,7 @@ bool device_impl::has(aspect Aspect) const {
                           BE == sycl::backend::opencl;
     return (is_cpu() || is_gpu()) && isCompatibleBE;
   }
-  case aspect::ext_intel_2d_block_load: {
+  case aspect::ext_intel_esimd_2d_block_io: {
     ur_exp_device_2d_block_array_capability_flags_t BlockArrayCapabilities;
     bool CallSuccessful =
         getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
@@ -789,21 +789,8 @@ bool device_impl::has(aspect Aspect) const {
     }
 
     return BlockArrayCapabilities &
-           UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD;
-  }
-  case aspect::ext_intel_2d_block_store: {
-    ur_exp_device_2d_block_array_capability_flags_t BlockArrayCapabilities;
-    bool CallSuccessful =
-        getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
-            MDevice, UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP,
-            sizeof(BlockArrayCapabilities), &BlockArrayCapabilities,
-            nullptr) == UR_RESULT_SUCCESS;
-    if (!CallSuccessful) {
-      return false;
-    }
-
-    return BlockArrayCapabilities &
-           UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE;
+           (UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD |
+            UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE);
   }
   }
 

sycl/test-e2e/Basic/aspects.cpp

@@ -93,11 +93,8 @@ int main() {
     if (plt.has(aspect::ext_oneapi_virtual_functions)) {
       std::cout << "  ext_oneapi_virtual_functions" << std::endl;
     }
-    if (plt.has(aspect::ext_intel_2d_block_load)) {
-      std::cout << "  ext_intel_2d_block_load" << std::endl;
-    }
-    if (plt.has(aspect::ext_intel_2d_block_store)) {
-      std::cout << "  ext_intel_2d_block_store" << std::endl;
+    if (plt.has(aspect::ext_intel_esimd_2d_block_io)) {
+      std::cout << "  ext_intel_esimd_2d_block_io" << std::endl;
     }
   }
   std::cout << "Passed." << std::endl;

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

@@ -223,14 +223,13 @@ TEST_F(DeviceInfoTest, BuiltInKernelIDs) {
       msg, "Attempting to use a built-in kernel. They are not fully supported");
 }
 
-TEST_F(DeviceInfoTest, GetDevice2DBlockArray) {
+TEST_F(DeviceInfoTest, GetDeviceESIMD2DBlockIO) {
   context Ctx{Plt.get_devices()[0]};
   TestContext.reset(new TestCtx(Ctx));
 
   device Dev = Ctx.get_devices()[0];
 
-  EXPECT_TRUE(Dev.has(aspect::ext_intel_2d_block_load));
-  EXPECT_TRUE(Dev.has(aspect::ext_intel_2d_block_store));
+  EXPECT_TRUE(Dev.has(aspect::ext_intel_esimd_2d_block_io));
 }
 
 TEST_F(DeviceInfoNegativeTest, TestAspectNotSupported) {
@@ -241,8 +240,7 @@ TEST_F(DeviceInfoNegativeTest, TestAspectNotSupported) {
   EXPECT_EQ(Dev.has(aspect::ext_intel_free_memory), false);
   EXPECT_EQ(Dev.has(aspect::ext_intel_memory_clock_rate), false);
   EXPECT_EQ(Dev.has(aspect::ext_intel_memory_bus_width), false);
-  EXPECT_EQ(Dev.has(aspect::ext_intel_2d_block_load), false);
-  EXPECT_EQ(Dev.has(aspect::ext_intel_2d_block_store), false);
+  EXPECT_EQ(Dev.has(aspect::ext_intel_esimd_2d_block_io), false);
 }
 
 TEST_F(DeviceInfoTest, SplitStringDelimeterSpace) {