From fbb25958951c74ceeb667d862b7dc4d244ad5e47 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 12 Jul 2023 16:12:18 -0700 Subject: [PATCH 1/5] Add support for local accessors to lsc API --- .../ext/intel/experimental/esimd/memory.hpp | 160 ++++++++++++++---- .../lsc_local_accessor_block_load_store.cpp | 114 +++++++++++++ .../lsc/lsc_local_accessor_gather_scatter.cpp | 121 +++++++++++++ 3 files changed, 366 insertions(+), 29 deletions(-) create mode 100644 sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp create mode 100644 sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp index e97ac2267eead..2a5e60bbaaf6c 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp @@ -718,8 +718,10 @@ template -__ESIMD_API std::enable_if_t, - __ESIMD_NS::simd> +__ESIMD_API std::enable_if_t< + !std::is_pointer_v && + !sycl::detail::acc_properties::is_local_accessor_v, + __ESIMD_NS::simd> lsc_gather(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM __ESIMD_NS::simd offsets, @@ -755,10 +757,12 @@ template -__ESIMD_API std::enable_if_t && - std::is_integral_v && - !std::is_same_v, - __ESIMD_NS::simd> +__ESIMD_API std::enable_if_t< + !std::is_pointer_v && + !sycl::detail::acc_properties::is_local_accessor_v + std::is_integral_v && + !std::is_same_v, + __ESIMD_NS::simd> lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, __ESIMD_NS::simd_mask pred = 1) { return lsc_gather( @@ -766,6 +770,19 @@ lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, } #endif +template +__ESIMD_API std::enable_if_t< + sycl::detail::acc_properties::is_local_accessor_v, + __ESIMD_NS::simd> +lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, + __ESIMD_NS::simd_mask pred = 1) { + return lsc_slm_gather( + offsets + __ESIMD_DNS::localAccessorToOffset(acc), pred); +} + /// Accessor-based gather. /// Supported platforms: DG2, PVC /// VISA instruction: lsc_load.ugm @@ -791,8 +808,10 @@ template -__ESIMD_API std::enable_if_t, - __ESIMD_NS::simd> +__ESIMD_API std::enable_if_t< + !std::is_pointer_v && + !sycl::detail::acc_properties::is_local_accessor_v, + __ESIMD_NS::simd> lsc_gather(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM __ESIMD_NS::simd offsets, @@ -832,10 +851,11 @@ template -__ESIMD_API std::enable_if_t && - std::is_integral_v && - !std::is_same_v, - __ESIMD_NS::simd> +__ESIMD_API std::enable_if_t< + !std::is_pointer_v && + !sycl::detail::acc_properties::is_local_accessor_v && + std::is_integral_v && !std::is_same_v, + __ESIMD_NS::simd> lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, __ESIMD_NS::simd_mask pred, __ESIMD_NS::simd old_values) { @@ -844,6 +864,20 @@ lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, } #endif +template +__ESIMD_API std::enable_if_t< + sycl::detail::acc_properties::is_local_accessor_v, + __ESIMD_NS::simd> +lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, + __ESIMD_NS::simd_mask pred, + __ESIMD_NS::simd old_values) { + return lsc_slm_gather( + offsets + __ESIMD_DNS::localAccessorToOffset(acc), pred, old_values); +} + /// USM pointer transposed gather with 1 channel. /// Supported platforms: DG2, PVC /// VISA instruction: lsc_load.ugm @@ -1135,9 +1169,11 @@ template -__ESIMD_API std::enable_if_t::value && - __ESIMD_NS::is_simd_flag_type_v, - __ESIMD_NS::simd> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v && + __ESIMD_NS::is_simd_flag_type_v, + __ESIMD_NS::simd> lsc_block_load(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM uint64_t offset, @@ -1208,6 +1244,20 @@ lsc_block_load(AccessorTy acc, #endif // !__ESIMD_FORCE_STATELESS_MEM } +template +__ESIMD_API std::enable_if_t< + sycl::detail::acc_properties::is_local_accessor_v && + __ESIMD_NS::is_simd_flag_type_v, + __ESIMD_NS::simd> +lsc_block_load(AccessorTy acc, uint32_t offset, + __ESIMD_NS::simd_mask<1> pred = 1, FlagsT flags = FlagsT{}) { + return lsc_slm_block_load( + offset + __ESIMD_DNS::localAccessorToOffset(acc), pred); +} + /// A variation of lsc_block_load without predicate parameter to simplify use /// of alignment parameter /// @@ -1299,9 +1349,11 @@ template -__ESIMD_API std::enable_if_t::value && - __ESIMD_NS::is_simd_flag_type_v, - __ESIMD_NS::simd> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v && + __ESIMD_NS::is_simd_flag_type_v, + __ESIMD_NS::simd> lsc_block_load(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM uint64_t offset, @@ -1372,6 +1424,19 @@ lsc_block_load(AccessorTy acc, #endif // !__ESIMD_FORCE_STATELESS_MEM } +template +__ESIMD_API std::enable_if_t< + sycl::detail::acc_properties::is_local_accessor_v && + __ESIMD_NS::is_simd_flag_type_v, + __ESIMD_NS::simd> +lsc_block_load(AccessorTy acc, uint32_t offset, __ESIMD_NS::simd_mask<1> pred, + __ESIMD_NS::simd old_values, FlagsT flags = FlagsT{}) { + return lsc_slm_block_load( + offset + __ESIMD_DNS::localAccessorToOffset(acc), pred, old_values); +} /// USM pointer prefetch gather. /// Supported platforms: DG2, PVC /// VISA instruction: lsc_load.ugm @@ -1493,7 +1558,9 @@ template -__ESIMD_API std::enable_if_t::value> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v> lsc_prefetch(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM __ESIMD_NS::simd offsets, @@ -1527,9 +1594,10 @@ template -__ESIMD_API std::enable_if_t::value && - std::is_integral_v && - !std::is_same_v> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v && + std::is_integral_v && !std::is_same_v> lsc_prefetch(AccessorTy acc, __ESIMD_NS::simd offsets, __ESIMD_NS::simd_mask pred = 1) { lsc_prefetch( @@ -1556,7 +1624,9 @@ template -__ESIMD_API std::enable_if_t::value> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v> lsc_prefetch(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM uint64_t offset @@ -1747,7 +1817,9 @@ template -__ESIMD_API std::enable_if_t::value> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v> lsc_scatter(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM __ESIMD_NS::simd offsets, @@ -1785,9 +1857,10 @@ template -__ESIMD_API std::enable_if_t::value && - std::is_integral_v && - !std::is_same_v> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v && + std::is_integral_v && !std::is_same_v> lsc_scatter(AccessorTy acc, __ESIMD_NS::simd offsets, __ESIMD_NS::simd vals, __ESIMD_NS::simd_mask pred = 1) { @@ -1795,6 +1868,20 @@ lsc_scatter(AccessorTy acc, __ESIMD_NS::simd offsets, acc, convert(offsets), vals, pred); } #endif + +template +__ESIMD_API std::enable_if_t< + sycl::detail::acc_properties::is_local_accessor_v> +lsc_scatter(AccessorTy acc, __ESIMD_NS::simd offsets, + __ESIMD_NS::simd vals, + __ESIMD_NS::simd_mask pred = 1) { + lsc_slm_scatter( + offsets + __ESIMD_DNS::localAccessorToOffset(acc), vals, pred); +} + /// USM pointer transposed scatter with 1 channel. /// Supported platforms: DG2, PVC /// VISA instruction: lsc_store.ugm @@ -1966,8 +2053,10 @@ template -__ESIMD_API std::enable_if_t::value && - __ESIMD_NS::is_simd_flag_type_v> +__ESIMD_API std::enable_if_t< + !std::is_pointer::value && + !sycl::detail::acc_properties::is_local_accessor_v && + __ESIMD_NS::is_simd_flag_type_v> lsc_block_store(AccessorTy acc, #ifdef __ESIMD_FORCE_STATELESS_MEM uint64_t offset, @@ -2041,6 +2130,19 @@ lsc_block_store(AccessorTy acc, #endif } +template +__ESIMD_API std::enable_if_t< + sycl::detail::acc_properties::is_local_accessor_v && + __ESIMD_NS::is_simd_flag_type_v> +lsc_block_store(AccessorTy acc, uint32_t offset, + __ESIMD_NS::simd vals, FlagsT flags = FlagsT{}) { + lsc_slm_block_store( + offset + __ESIMD_DNS::localAccessorToOffset(acc), vals); +} + /// A variation of lsc_block_store without predicate parameter to simplify /// use of alignment parameter /// diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp new file mode 100644 index 0000000000000..612a4a86028d2 --- /dev/null +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp @@ -0,0 +1,114 @@ +//==- lsc_local_accessor_block_load_store.cpp - DPC++ ESIMD on-device test ==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu-intel-pvc +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// +// TODO: Enable the test when GPU driver is ready/fixed. +// XFAIL: opencl || windows || gpu-intel-pvc +// TODO: add support for local_accessors to esimd_emulator. +// UNSUPPORTED: esimd_emulator +// This test verifies usage of block_load/block_store for local_accessor. + +#include "../esimd_test_utils.hpp" + +#include +#include + +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; +using namespace sycl::ext::intel::experimental::esimd; + +constexpr int VL = 16; + +template +bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) { + std::cout << "Running case: T=" << esimd_test::type_name() << std::endl; + + // The test is going to use (LocalRange * VL) elements of T type. + auto Dev = Q.get_device(); + auto DeviceSLMSize = Dev.get_info(); + if (DeviceSLMSize < LocalRange * VL * sizeof(T)) { + // Report an error - the test needs a fix. + std::cerr << "Error: Test needs more SLM memory than device has!" + << std::endl; + return false; + } + + T *Out = malloc_shared(GlobalRange * VL, Q); + for (int I = 0; I < GlobalRange * VL; I++) + Out[I] = -1; + + try { + nd_range<1> NDRange{range<1>{GlobalRange}, range<1>{LocalRange}}; + Q.submit([&](handler &CGH) { + auto LocalAcc = local_accessor(LocalRange * VL, CGH); + + CGH.parallel_for(NDRange, [=](nd_item<1> Item) SYCL_ESIMD_KERNEL { + uint32_t GID = Item.get_global_id(0); + uint32_t LID = Item.get_local_id(0); + + simd IntValues(GID * 100, 1); + simd ValuesToSLM = IntValues; + lsc_block_store(LocalAcc, LID * VL * sizeof(T), ValuesToSLM); + + Item.barrier(); + + if (LID == 0) { + for (int LID = 0; LID < LocalRange; LID++) { + simd ValuesFromSLM = + lsc_block_load(LocalAcc, LID * VL * sizeof(T)); + ValuesFromSLM.copy_to(Out + (GID + LID) * VL); + } // end for (int LID = 0; LID < LocalRange; LID++) + } // end if (LID == 0) + }); + }).wait(); + } catch (sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + free(Out, Q); + return false; + } + + bool Pass = true; + for (int I = 0; I < GlobalRange * VL; I++) { + int GID = I / VL; + int LID = GID % LocalRange; + int VecElementIndex = I % VL; + + T Expected = GID * 100 + VecElementIndex; + T Computed = Out[I]; + if (Computed != Expected) { + std::cout << "Error: Out[" << I << "]:" << Computed << " != " << Expected + << ":[expected]" << std::endl; + Pass = false; + } + } + + free(Out, Q); + return Pass; +} + +int main() { + auto Q = queue{gpu_selector_v}; + auto Dev = Q.get_device(); + auto DeviceSLMSize = Dev.get_info(); + std::cout << "Running on " << Dev.get_info() + << ", Local memory size available : " << DeviceSLMSize << std::endl; + + constexpr uint32_t LocalRange = 16; + constexpr uint32_t GlobalRange = LocalRange * 2; // 2 groups. + + bool Pass = true; + Pass &= test(Q, LocalRange, GlobalRange); + Pass &= test(Q, LocalRange, GlobalRange); + + std::cout << "Test result: " << (Pass ? "Pass" : "Fail") << std::endl; + return Pass ? 0 : 1; +} diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp new file mode 100644 index 0000000000000..d92a9120fac2f --- /dev/null +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp @@ -0,0 +1,121 @@ +//==-- lsc_local_accessor_gather_scatter.cpp - DPC++ ESIMD on-device test -==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu-intel-pvc +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// +// TODO: Enable the test when GPU driver is ready/fixed. +// XFAIL: opencl || windows || gpu-intel-pvc +// TODO: add support for local_accessors to esimd_emulator. +// UNSUPPORTED: esimd_emulator +// The test checks functionality of the gather/scatter local +// accessor-based ESIMD intrinsics. + +#include "../esimd_test_utils.hpp" + +#include +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::experimental::esimd; + +template bool test(queue q) { + constexpr size_t size = VL; + constexpr int MASKED_LANE = VL - 1; + + std::cout << "Testing T=" << typeid(T).name() << " VL=" << VL + << " STRIDE=" << STRIDE << "...\n"; + + // The test is going to use size elements of T type. + auto Dev = q.get_device(); + auto DeviceSLMSize = Dev.get_info(); + if (DeviceSLMSize < size) { + // Report an error - the test needs a fix. + std::cerr << "Error: Test needs more SLM memory than device has!" + << std::endl; + return false; + } + + T *A = new T[size]; + + for (unsigned i = 0; i < size; ++i) { + A[i] = static_cast(i); + } + + try { + buffer buf(A, range<1>(size)); + nd_range<1> glob_range{range<1>{1}, range<1>{1}}; + + q.submit([&](handler &cgh) { + auto acc = buf.template get_access(cgh); + auto LocalAcc = local_accessor(size * STRIDE, cgh); + cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::ext::intel::esimd; + simd valsIn; + valsIn.copy_from(acc, 0); + simd_mask pred = 1; + pred[MASKED_LANE] = 0; // mask out the last lane + LocalAcc[MASKED_LANE * STRIDE] = -1; + simd offsets(0, STRIDE * sizeof(T)); + lsc_scatter(LocalAcc, offsets, valsIn, pred); + + simd valsOut = lsc_gather(LocalAcc, offsets); + + valsOut.copy_to(acc, 0); + }); + }).wait(); + } catch (sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + delete[] A; + return false; // not success + } + + int err_cnt = 0; + + for (unsigned i = 0; i < size; ++i) { + T gold = i != MASKED_LANE ? static_cast(i) : static_cast(-1); + + if (A[i] != gold) { + if (++err_cnt < 35) { + std::cout << "failed at index " << i << ": " << A[i] << " != " << gold + << " (gold)\n"; + } + } + } + if (err_cnt > 0) { + std::cout << " pass rate: " + << ((float)(size - err_cnt) / (float)size) * 100.0f << "% (" + << (size - err_cnt) << "/" << size << ")\n"; + } + + delete[] A; + + std::cout << (err_cnt > 0 ? " FAILED\n" : " Passed\n"); + return err_cnt > 0 ? false : true; +} + +int main(void) { + queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() << "\n"; + + bool passed = true; + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + passed &= test(q); + return passed ? 0 : 1; +} From 1d1176bdef9bca5afaae8b16f0e17f1a40130d2f Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 12 Jul 2023 16:37:27 -0700 Subject: [PATCH 2/5] Fix clang-format issue --- sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp index 612a4a86028d2..80cde87452dbd 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp @@ -108,7 +108,7 @@ int main() { bool Pass = true; Pass &= test(Q, LocalRange, GlobalRange); Pass &= test(Q, LocalRange, GlobalRange); - + std::cout << "Test result: " << (Pass ? "Pass" : "Fail") << std::endl; return Pass ? 0 : 1; } From d6caa60d26bec5c3e4312281099e633d8a0a2ac1 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 12 Jul 2023 19:13:22 -0700 Subject: [PATCH 3/5] Fix a typo --- sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp index 2a5e60bbaaf6c..d56e3f37603e5 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp @@ -759,8 +759,8 @@ template __ESIMD_API std::enable_if_t< !std::is_pointer_v && - !sycl::detail::acc_properties::is_local_accessor_v - std::is_integral_v && + !sycl::detail::acc_properties::is_local_accessor_v && + std::is_integral_v && !std::is_same_v, __ESIMD_NS::simd> lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, From 3a98a00e984c089a56420326bce53ebc2b0b39be Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 12 Jul 2023 19:17:42 -0700 Subject: [PATCH 4/5] Fix clang-format issue --- sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp index d56e3f37603e5..6f73ef7f55e2c 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp +++ b/sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp @@ -760,8 +760,7 @@ template && !sycl::detail::acc_properties::is_local_accessor_v && - std::is_integral_v && - !std::is_same_v, + std::is_integral_v && !std::is_same_v, __ESIMD_NS::simd> lsc_gather(AccessorTy acc, __ESIMD_NS::simd offsets, __ESIMD_NS::simd_mask pred = 1) { From 1629b5b5b6078b88f9d312e4e68efb4eb8e92dcc Mon Sep 17 00:00:00 2001 From: gregory Date: Mon, 24 Jul 2023 17:54:07 -0700 Subject: [PATCH 5/5] Update test comments to address PR comments --- sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp | 1 - sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp | 1 - 2 files changed, 2 deletions(-) diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp index 80cde87452dbd..31af2d96a677e 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp @@ -11,7 +11,6 @@ // // TODO: Enable the test when GPU driver is ready/fixed. // XFAIL: opencl || windows || gpu-intel-pvc -// TODO: add support for local_accessors to esimd_emulator. // UNSUPPORTED: esimd_emulator // This test verifies usage of block_load/block_store for local_accessor. diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp index d92a9120fac2f..31307a9203259 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_gather_scatter.cpp @@ -11,7 +11,6 @@ // // TODO: Enable the test when GPU driver is ready/fixed. // XFAIL: opencl || windows || gpu-intel-pvc -// TODO: add support for local_accessors to esimd_emulator. // UNSUPPORTED: esimd_emulator // The test checks functionality of the gather/scatter local // accessor-based ESIMD intrinsics.