diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 72841724fa01d..e9d52c5b23dac 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/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") diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index 8fff99790470a..b5c82c3ead941 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -1,7 +1,7 @@ -# commit 0a90db9b2c36960c9b28ce18557ca15760724c4d -# Merge: c4d9fdb4 6e0bdeb9 +# commit db83117e830406b0d9950e24892dba868acba354 +# Merge: 0a90db9b c79df596 # Author: Callum Fare -# 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) diff --git a/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md b/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md index c01223a2e4fe0..6319e9e4751e2 100644 --- a/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md +++ b/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 diff --git a/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md b/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md index 5fb775dddc3bb..ae897ebc59d5a 100644 --- a/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md +++ b/sycl/doc/extensions/supported/sycl_ext_intel_esimd/sycl_ext_intel_esimd_functions.md @@ -552,7 +552,7 @@ Loads and returns a vector `simd` 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` 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. diff --git a/sycl/include/sycl/info/ext_intel_device_traits.def b/sycl/include/sycl/info/ext_intel_device_traits.def index f828b43e2a1d7..50b4e9eec952a 100644 --- a/sycl/include/sycl/info/ext_intel_device_traits.def +++ b/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 diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index 2862913194ae0..149d3f96eb0a5 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -1589,6 +1589,25 @@ get_device_info( return get_device_info_impl::get(Dev); } +template <> +inline ext::intel::esimd::info::device::has_2d_block_io_support::return_type +get_device_info( + 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( + 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 diff --git a/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp b/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp index 03fa6a39cea13..88285426ef900 100644 --- a/sycl/test-e2e/ESIMD/InlineAsm/asm_glb.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; auto e = q.submit([&](handler &cgh) { auto PA = bufa.get_access(cgh); diff --git a/sycl/test-e2e/ESIMD/addc.cpp b/sycl/test-e2e/ESIMD/addc.cpp index 30a73d900ea3f..956b860e97b53 100644 --- a/sycl/test-e2e/ESIMD/addc.cpp +++ b/sycl/test-e2e/ESIMD/addc.cpp @@ -138,7 +138,7 @@ template 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() << "\n"; + std::cout << "Running on " << D.get_info() << "\n"; constexpr bool AIsVector = true; constexpr bool BIsVector = true; diff --git a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp b/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp index 999ab2f91bfec..11a3a02aede09 100644 --- a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; auto e = q.submit([&](handler &cgh) { auto PA = bufa.get_access(cgh); diff --git a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp b/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp index 8b2b46e046ee7..4537ddd210cb9 100644 --- a/sycl/test-e2e/ESIMD/api/functional/operators/operator_assignment_glb_mask.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; auto e = q.submit([&](handler &cgh) { auto PA = bufa.get_access(cgh); diff --git a/sycl/test-e2e/ESIMD/dpas/dpas_bf16.cpp b/sycl/test-e2e/ESIMD/dpas/dpas_bf16.cpp index 876fe7c9466cf..578cca3eff771 100644 --- a/sycl/test-e2e/ESIMD/dpas/dpas_bf16.cpp +++ b/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() << std::endl; + std::cout << "Running on " << Dev.get_info() + << std::endl; bool Print = argc > 1 && std::string(argv[1]) == "-debug"; bool Passed = true; diff --git a/sycl/test-e2e/ESIMD/dpas/dpas_fp16.cpp b/sycl/test-e2e/ESIMD/dpas/dpas_fp16.cpp index 278f8397018d3..601dc107ed9f8 100644 --- a/sycl/test-e2e/ESIMD/dpas/dpas_fp16.cpp +++ b/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() << std::endl; + std::cout << "Running on " << Dev.get_info() + << std::endl; bool Print = argc > 1 && std::string(argv[1]) == "-debug"; bool Passed = true; diff --git a/sycl/test-e2e/ESIMD/dpas/dpas_int.cpp b/sycl/test-e2e/ESIMD/dpas/dpas_int.cpp index 80cd9a0f4be5f..6156456d1b387 100644 --- a/sycl/test-e2e/ESIMD/dpas/dpas_int.cpp +++ b/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() << std::endl; + std::cout << "Running on " << Dev.get_info() + << std::endl; bool Print = argc > 1 && std::string(argv[1]) == "-debug"; bool Passed = true; diff --git a/sycl/test-e2e/ESIMD/imulh_umulh.cpp b/sycl/test-e2e/ESIMD/imulh_umulh.cpp index 9fba387948f07..fe74bf72d80a9 100644 --- a/sycl/test-e2e/ESIMD/imulh_umulh.cpp +++ b/sycl/test-e2e/ESIMD/imulh_umulh.cpp @@ -155,7 +155,7 @@ template 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() << "\n"; + std::cout << "Running on " << D.get_info() << "\n"; constexpr bool AIsVector = true; constexpr bool BIsVector = true; diff --git a/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp b/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp index 59c50e0547bfb..52318cb483688 100644 --- a/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; bool passed = true; passed &= test(q); diff --git a/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp b/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp index 6fcbcef600557..78f40e4377be3 100644 --- a/sycl/test-e2e/ESIMD/lsc/atomic_smoke.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; Config cfg{ 11, // int threads_per_group; diff --git a/sycl/test-e2e/ESIMD/lsc/local_accessor_atomic_smoke.cpp b/sycl/test-e2e/ESIMD/lsc/local_accessor_atomic_smoke.cpp index bb54ca891b4c7..e2eb04d883db3 100644 --- a/sycl/test-e2e/ESIMD/lsc/local_accessor_atomic_smoke.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; bool passed = true; #ifndef CMPXCHG_TEST diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp index 8067859151762..c289b7a61da92 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_argument_type_deduction.cpp +++ b/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() + std::cout << "Device name: " << device.get_info() << std::endl; int error = testUSM<8>(q); diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp index f1e124dcf8a7e..af91f470002a1 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_predicate.cpp +++ b/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() + std::cout << "Device name: " << device.get_info() << std::endl; int error = testUSM<8>(q); diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp index 2a978721e070b..eba3f4cba5663 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_predicate_stateless.cpp +++ b/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() + std::cout << "Device name: " << device.get_info() << std::endl; int error = testAccessor<8>(q); diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_smoke.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_smoke.cpp index 85b052eaf5f3a..93db1fef805ac 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_smoke.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; bool passed = true; #ifndef CMPXCHG_TEST diff --git a/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp b/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp index d98490c871eef..c3b9282d9bdf0 100644 --- a/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp +++ b/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp @@ -149,8 +149,8 @@ template bool tests(queue Q) { int main() { queue Q; - std::cout << "Running on " << Q.get_device().get_info() - << "\n"; + std::cout << "Running on " + << Q.get_device().get_info() << "\n"; bool Passed = true; Passed &= tests(Q); diff --git a/sycl/test-e2e/ESIMD/radix_sort.cpp b/sycl/test-e2e/ESIMD/radix_sort.cpp index ca62ff61a88b8..bfc3dfa53ae00 100644 --- a/sycl/test-e2e/ESIMD/radix_sort.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; auto ctxt = q.get_context(); // allocate and initialized input diff --git a/sycl/test-e2e/ESIMD/slm_alloc.cpp b/sycl/test-e2e/ESIMD/slm_alloc.cpp index ef5d191ad56d7..e0c4f8c102151 100644 --- a/sycl/test-e2e/ESIMD/slm_alloc.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; std::cout << "force_inline=" << force_inline << "\n"; auto ctxt = q.get_context(); uint32_t size = SLM_TOTAL * NUM_WGS / ELEM_SIZE; diff --git a/sycl/test-e2e/ESIMD/slm_alloc_many_kernels_many_funcs.cpp b/sycl/test-e2e/ESIMD/slm_alloc_many_kernels_many_funcs.cpp index c58423d747ecd..a00d77e338ea7 100644 --- a/sycl/test-e2e/ESIMD/slm_alloc_many_kernels_many_funcs.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; std::cout << "force_inline=" << force_inline << "\n"; auto ctxt = q.get_context(); diff --git a/sycl/test-e2e/ESIMD/slm_alloc_many_kernels_one_func.cpp b/sycl/test-e2e/ESIMD/slm_alloc_many_kernels_one_func.cpp index 818990eb2d13f..3a9dea2949c8e 100644 --- a/sycl/test-e2e/ESIMD/slm_alloc_many_kernels_one_func.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; std::cout << "force_inline=" << force_inline << "\n"; auto ctxt = q.get_context(); diff --git a/sycl/test-e2e/ESIMD/subb.cpp b/sycl/test-e2e/ESIMD/subb.cpp index d577be1662860..0174501b6cf02 100644 --- a/sycl/test-e2e/ESIMD/subb.cpp +++ b/sycl/test-e2e/ESIMD/subb.cpp @@ -140,7 +140,7 @@ template 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() << "\n"; + std::cout << "Running on " << D.get_info() << "\n"; constexpr bool AIsVector = true; constexpr bool BIsVector = true; diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm.cpp index 430b834efe8a7..52aa751c887e7 100644 --- a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; bool passed = true; diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp index cc38560f5de57..281727c610184 100644 --- a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; bool passed = true; diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/device_info_descriptors.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/device_info_descriptors.cpp new file mode 100644 index 0000000000000..f7e85a9407693 --- /dev/null +++ b/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 +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + auto Arch = Q.get_device().get_info(); + 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; +} diff --git a/sycl/test-e2e/ESIMD/vadd_raw_send_gen12.cpp b/sycl/test-e2e/ESIMD/vadd_raw_send_gen12.cpp index 7e5d8fff00e8e..d46907dae5bdc 100644 --- a/sycl/test-e2e/ESIMD/vadd_raw_send_gen12.cpp +++ b/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() << "\n"; + std::cout << "Running on " << dev.get_info() + << "\n"; int *A = malloc_shared(Size, q); int *B = malloc_shared(Size, q); diff --git a/sycl/test-e2e/ESIMD/wait.cpp b/sycl/test-e2e/ESIMD/wait.cpp index 0f37ab61caae0..df0c2dfb955c9 100644 --- a/sycl/test-e2e/ESIMD/wait.cpp +++ b/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() << std::endl; + std::cout << "Running on " << Dev.get_info() + << std::endl; bool Passed = true; Passed &= test(Q); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 853ac28bad1d5..983a69b2479aa 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3677,6 +3677,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 @@ -3782,6 +3783,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 diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 8e29aba2726c9..db76840dfad3b 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -67,6 +67,7 @@ ??$get_info@Ugpu_hw_threads_per_eu@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ ??$get_info@Ugpu_slices@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ ??$get_info@Ugpu_subslices_per_slice@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ +??$get_info@Uhas_2d_block_io_support@device@info@esimd@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA_NXZ ??$get_info@Uimage_row_pitch_align@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ ??$get_info@Umatrix_combinations@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@Umax_compute_queue_indices@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAHXZ @@ -156,6 +157,7 @@ ??$get_info_impl@Ugpu_slices@device@info@intel@ext@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ ??$get_info_impl@Ugpu_subslices_per_slice@device@info@intel@ext@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ ??$get_info_impl@Uhalf_fp_config@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4fp_config@info@_V1@sycl@@V?$allocator@W4fp_config@info@_V1@sycl@@@std@@@std@@XZ +??$get_info_impl@Uhas_2d_block_io_support@device@info@esimd@intel@ext@_V1@sycl@@@device@_V1@sycl@@AEBA_NXZ ??$get_info_impl@Uhost_unified_memory@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_NXZ ??$get_info_impl@Uimage2d_max_height@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_KXZ ??$get_info_impl@Uimage2d_max_width@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA_KXZ diff --git a/sycl/unittests/kernel-and-program/DeviceInfo.cpp b/sycl/unittests/kernel-and-program/DeviceInfo.cpp index 532d656ac49c1..3c7f06b75cb81 100644 --- a/sycl/unittests/kernel-and-program/DeviceInfo.cpp +++ b/sycl/unittests/kernel-and-program/DeviceInfo.cpp @@ -27,6 +27,9 @@ struct TestCtx { static std::unique_ptr TestContext; +ur_exp_device_2d_block_array_capability_flags_t Flags2DBlockIO = 0; +bool HasESIMDSupport = false; + static ur_result_t redefinedDeviceGetInfo(void *pParams) { auto params = *static_cast(pParams); if (*params.ppropName == UR_DEVICE_INFO_UUID) { @@ -72,6 +75,20 @@ 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( + *params.ppPropValue) = Flags2DBlockIO; + } + } + + if (*params.ppropName == UR_DEVICE_INFO_ESIMD_SUPPORT) { + assert(*params.ppropSize == sizeof(bool)); + if (*params.ppPropValue) + *static_cast(*params.ppPropValue) = HasESIMDSupport; + } return UR_RESULT_SUCCESS; } @@ -185,6 +202,42 @@ 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]; + + HasESIMDSupport = true; + Flags2DBlockIO = UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD | + UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE; + auto HasSupport = + Dev.get_info(); + EXPECT_TRUE(HasSupport); + + Flags2DBlockIO = UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD; + HasSupport = + Dev.get_info(); + EXPECT_FALSE(HasSupport); + + Flags2DBlockIO = UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE; + HasSupport = + Dev.get_info(); + EXPECT_FALSE(HasSupport); + + Flags2DBlockIO = 0; + HasSupport = + Dev.get_info(); + EXPECT_FALSE(HasSupport); + + Flags2DBlockIO = UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD | + UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE; + HasESIMDSupport = false; + HasSupport = + Dev.get_info(); + EXPECT_FALSE(HasSupport); +} + TEST_F(DeviceInfoTest, BuiltInKernelIDs) { context Ctx{Plt.get_devices()[0]}; TestContext.reset(new TestCtx(Ctx));