Skip to content

Commit 474461c

Browse files
fineg74fineg74
authored
[SYCL][ESIMD] Add support for local accessors to copy_from/copy_to API (intel#10310)
1 parent bf128c8 commit 474461c

File tree

8 files changed

+157
-23
lines changed

8 files changed

+157
-23
lines changed

‎sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp

Lines changed: 7 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -324,12 +324,13 @@ class simd_obj_impl {
324324
/// argument.
325325
/// @param acc The accessor to read from.
326326
/// @param offset offset in bytes of the first element.
327-
template <
328-
typename AccessorT, typename Flags = element_aligned_tag,
329-
typename = std::enable_if_t<
330-
detail::is_sycl_accessor_with<AccessorT, accessor_mode_cap::can_read,
331-
sycl::access::target::device>::value &&
332-
is_simd_flag_type_v<Flags>>>
327+
template <typename AccessorT, typename Flags = element_aligned_tag,
328+
typename = std::enable_if_t<
329+
(sycl::detail::acc_properties::is_local_accessor_v<AccessorT> ||
330+
detail::is_sycl_accessor_with<
331+
AccessorT, accessor_mode_cap::can_read,
332+
sycl::access::target::device>::value) &&
333+
is_simd_flag_type_v<Flags>>>
333334
simd_obj_impl(AccessorT acc,
334335
#ifdef __ESIMD_FORCE_STATELESS_MEM
335336
uint64_t offset,

‎sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -77,7 +77,9 @@ struct is_sycl_accessor_with
7777
template <typename T, accessor_mode_cap_val_t Capability,
7878
sycl::access::target AccessTarget, typename RetT>
7979
using EnableIfAccessor = std::enable_if_t<
80-
detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value, RetT>;
80+
detail::is_sycl_accessor_with<T, Capability, AccessTarget>::value ||
81+
sycl::detail::acc_properties::is_local_accessor_v<T>,
82+
RetT>;
8183

8284
template <typename T, int Dimensions>
8385
__ESIMD_API uint32_t localAccessorToOffset(local_accessor<T, Dimensions> acc) {

‎sycl/include/sycl/ext/intel/esimd/memory.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2019,6 +2019,7 @@ block_load(AccessorTy acc, uint32_t offset, Flags = {}) {
20192019
__ESIMD_DNS::localAccessorToOffset(acc));
20202020
}
20212021

2022+
/// Variant of block_store that uses local accessor as a parameter.
20222023
/// Stores elements of the vector \p vals to a contiguous block of SLM memory
20232024
/// represented by the given local accessor and the byte-offset \p offset.
20242025
/// The generated code depends on the combination {T, N, Flags}.

‎sycl/test-e2e/ESIMD/local_accessor_block_load_store.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,9 +11,6 @@
1111
// https://github.com/intel/llvm/issues/10369
1212
// UNSUPPORTED: gpu
1313
//
14-
// TODO: Enable the test when GPU driver is ready/fixed.
15-
// XFAIL: opencl || windows || gpu-intel-pvc
16-
// TODO: add support for local_accessors to esimd_emulator.
1714
// UNSUPPORTED: esimd_emulator
1815
// This test verifies usage of block_load/block_store for local_accessor.
1916

‎sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp

Lines changed: 140 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,140 @@
1+
//==-- local_accessor_copy_to_from.cpp - DPC++ ESIMD on-device test -------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// RUN: %{build} -o %t.out
9+
// RUN: %{run} %t.out
10+
//
11+
// UNSUPPORTED: gpu
12+
// UNSUPPORTED: esimd_emulator
13+
// The test checks functionality of the gather/scatter local
14+
// accessor-based ESIMD intrinsics.
15+
16+
#include "esimd_test_utils.hpp"
17+
18+
#include <sycl/ext/intel/esimd.hpp>
19+
#include <sycl/sycl.hpp>
20+
21+
#include <iostream>
22+
23+
using namespace sycl;
24+
using namespace sycl::ext::intel::esimd;
25+
26+
constexpr uint32_t LocalRange = 16;
27+
constexpr uint32_t GlobalRange = LocalRange * 2; // 2 groups.
28+
29+
template <typename T, unsigned VL> bool test(queue q) {
30+
constexpr size_t Size = VL * LocalRange;
31+
std::cout << "Running case: T=" << esimd_test::type_name<T>() << " VL=" << VL
32+
<< std::endl;
33+
34+
// The test is going to use (LocalRange * VL) elements of T type.
35+
auto Dev = q.get_device();
36+
auto DeviceSLMSize = Dev.get_info<sycl::info::device::local_mem_size>();
37+
if (DeviceSLMSize < Size * sizeof(T)) {
38+
// Report an error - the test needs a fix.
39+
std::cerr << "Error: Test needs more SLM memory than device has!"
40+
<< std::endl;
41+
return false;
42+
}
43+
44+
T *A = new T[GlobalRange * VL];
45+
46+
for (unsigned i = 0; i < GlobalRange * VL; ++i) {
47+
A[i] = static_cast<T>(0);
48+
}
49+
50+
try {
51+
buffer<T, 1> buf(A, range<1>(GlobalRange * VL));
52+
nd_range<1> NDRange{range<1>{GlobalRange}, range<1>{LocalRange}};
53+
q.submit([&](handler &CGH) {
54+
auto LocalAcc = local_accessor<T, 1>(Size, CGH);
55+
auto Acc = buf.template get_access<access::mode::read_write>(CGH);
56+
CGH.parallel_for(NDRange, [=](nd_item<1> Item) SYCL_ESIMD_KERNEL {
57+
uint32_t GID = Item.get_global_id(0);
58+
uint32_t LID = Item.get_local_id(0);
59+
60+
simd<T, VL> ValuesToSLM(GID * 100, 1);
61+
ValuesToSLM.copy_to(LocalAcc, LID * VL * sizeof(T));
62+
63+
Item.barrier();
64+
65+
if (LID == 0) {
66+
for (int LID = 0; LID < LocalRange; LID++) {
67+
simd<T, VL> ValuesFromSLM;
68+
ValuesFromSLM.copy_from(LocalAcc, LID * VL * sizeof(T));
69+
ValuesFromSLM.copy_to(Acc, (GID + LID) * VL * sizeof(T));
70+
} // end for (int LID = 0; LID < LocalRange; LID++)
71+
} // end if (LID == 0)
72+
});
73+
}).wait();
74+
} catch (sycl::exception const &e) {
75+
std::cout << "SYCL exception caught: " << e.what() << '\n';
76+
delete[] A;
77+
return false;
78+
}
79+
80+
bool Pass = true;
81+
for (int I = 0; I < GlobalRange * VL; I++) {
82+
int GID = I / VL;
83+
int LID = GID % LocalRange;
84+
int VecElementIndex = I % VL;
85+
86+
T Expected = GID * 100 + VecElementIndex;
87+
T Computed = A[I];
88+
if (Computed != Expected) {
89+
std::cout << "Error: Out[" << I << "]:" << Computed << " != " << Expected
90+
<< ":[expected]" << std::endl;
91+
Pass = false;
92+
}
93+
}
94+
95+
delete[] A;
96+
97+
return Pass;
98+
}
99+
100+
int main() {
101+
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
102+
103+
auto dev = q.get_device();
104+
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
105+
106+
bool passed = true;
107+
passed &= test<char, 1>(q);
108+
passed &= test<char, 2>(q);
109+
passed &= test<char, 4>(q);
110+
passed &= test<char, 8>(q);
111+
passed &= test<char, 16>(q);
112+
passed &= test<char, 32>(q);
113+
passed &= test<char, 64>(q);
114+
passed &= test<char, 128>(q);
115+
passed &= test<short, 1>(q);
116+
passed &= test<short, 2>(q);
117+
passed &= test<short, 4>(q);
118+
passed &= test<short, 8>(q);
119+
passed &= test<short, 16>(q);
120+
passed &= test<short, 32>(q);
121+
passed &= test<short, 64>(q);
122+
passed &= test<short, 128>(q);
123+
passed &= test<int, 1>(q);
124+
passed &= test<int, 2>(q);
125+
passed &= test<int, 4>(q);
126+
passed &= test<int, 8>(q);
127+
passed &= test<int, 16>(q);
128+
passed &= test<int, 32>(q);
129+
passed &= test<int, 64>(q);
130+
passed &= test<int, 128>(q);
131+
passed &= test<float, 1>(q);
132+
passed &= test<float, 2>(q);
133+
passed &= test<float, 4>(q);
134+
passed &= test<float, 8>(q);
135+
passed &= test<float, 16>(q);
136+
passed &= test<float, 32>(q);
137+
passed &= test<float, 64>(q);
138+
passed &= test<float, 128>(q);
139+
return passed ? 0 : 1;
140+
}

‎sycl/test-e2e/ESIMD/local_accessor_gather_scatter.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,9 +11,6 @@
1111
// https://github.com/intel/llvm/issues/10369
1212
// UNSUPPORTED: gpu
1313
//
14-
// TODO: Enable the test when GPU driver is ready/fixed.
15-
// XFAIL: opencl || windows || gpu-intel-pvc
16-
// TODO: add support for local_accessors to esimd_emulator.
1714
// UNSUPPORTED: esimd_emulator
1815
// The test checks functionality of the gather/scatter local
1916
// accessor-based ESIMD intrinsics.

‎sycl/test-e2e/ESIMD/local_accessor_gather_scatter_rgba.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,9 +10,6 @@
1010
//
1111
// RUN: %{build} -o %t.out
1212
// RUN: %{run} %t.out
13-
// TODO: Enable the test when GPU driver is ready/fixed.
14-
// XFAIL: opencl || windows || gpu-intel-pvc
15-
// TODO: add support for local_accessors to esimd_emulator.
1613
// UNSUPPORTED: esimd_emulator
1714
// The test checks functionality of the gather_rgba/scatter_rgba local
1815
// accessor-based ESIMD intrinsics.

‎sycl/test/esimd/simd_copy_to_copy_from.cpp

Lines changed: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -34,25 +34,23 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
3434
v0.copy_to(ptr);
3535
}
3636

37-
// --- Negative tests.
38-
39-
// Incompatible target.
4037
SYCL_EXTERNAL void kernel3(local_accessor<int, 1> &buf) SYCL_ESIMD_FUNCTION {
4138
simd<int, 32> v1(0, 1);
4239
simd<int, 32> v0;
43-
// CHECK: simd_copy_to_copy_from.cpp:44{{.*}}error: no matching member function for call to 'copy_from'
4440
v0.copy_from(buf, 0);
4541
v0 = v0 + v1;
46-
// CHECK: simd_copy_to_copy_from.cpp:47{{.*}}error: no matching member function for call to 'copy_to'
4742
v0.copy_to(buf, 0);
4843
}
4944

45+
// --- Negative tests.
46+
5047
// Incompatible mode (write).
5148
SYCL_EXTERNAL void
5249
kernel4(accessor<int, 1, access::mode::write, access::target::device> &buf)
5350
SYCL_ESIMD_FUNCTION {
5451
simd<int, 32> v;
55-
// CHECK: simd_copy_to_copy_from.cpp:56{{.*}}error: no matching member function for call to 'copy_from'
52+
// CHECK: simd_copy_to_copy_from.cpp:54{{.*}}error: no matching member
53+
// function for call to 'copy_from'
5654
v.copy_from(buf, 0);
5755
}
5856

@@ -61,6 +59,7 @@ SYCL_EXTERNAL void
6159
kernel5(accessor<int, 1, access::mode::read, access::target::device> &buf)
6260
SYCL_ESIMD_FUNCTION {
6361
simd<int, 32> v(0, 1);
64-
// CHECK: simd_copy_to_copy_from.cpp:65{{.*}}error: no matching member function for call to 'copy_to'
62+
// CHECK: simd_copy_to_copy_from.cpp:64{{.*}}error: no matching member
63+
// function for call to 'copy_to'
6564
v.copy_to(buf, 0);
6665
}

0 commit comments

Comments
 (0)