Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit cb20604

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Add and adjust tests for SYCL 2020 multi_ptr (#1293)
This commit adjusts a number of existing using multi_ptr and adds new tests for new interfaces. Signed-off-by: Larsen, Steffen <[email protected]>
1 parent c7d9530 commit cb20604

File tree

6 files changed

+388
-110
lines changed

6 files changed

+388
-110
lines changed

‎SYCL/Basic/multi_ptr.cpp

Lines changed: 165 additions & 91 deletions
Original file line numberDiff line numberDiff line change
@@ -2,10 +2,6 @@
22
// RUN: %CPU_RUN_PLACEHOLDER %t.out
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
5-
// RUN: %clangxx -DRESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR -fsycl -fsycl-targets=%sycl_triple -fsycl-dead-args-optimization %s -o %t1.out
6-
// RUN: %CPU_RUN_PLACEHOLDER %t1.out
7-
// RUN: %GPU_RUN_PLACEHOLDER %t1.out
8-
// RUN: %ACC_RUN_PLACEHOLDER %t1.out
95

106
//==--------------- multi_ptr.cpp - SYCL multi_ptr test --------------------==//
117
//
@@ -24,8 +20,9 @@ using namespace sycl;
2420

2521
/* This is the class used to name the kernel for the runtime.
2622
* This must be done when the kernel is expressed as a lambda. */
27-
template <typename T> class testMultPtrKernel;
28-
template <typename T> class testMultPtrArrowOperatorKernel;
23+
template <typename T, access::decorated IsDecorated> class testMultPtrKernel;
24+
template <typename T, access::decorated IsDecorated>
25+
class testMultPtrArrowOperatorKernel;
2926

3027
template <typename T> struct point {
3128
point(const point &rhs) = default;
@@ -38,29 +35,60 @@ template <typename T> struct point {
3835
T y;
3936
};
4037

41-
template <typename T>
42-
void innerFunc(id<1> wiID, global_ptr<T> ptr_1, global_ptr<T> ptr_2,
43-
local_ptr<T> local_ptr) {
38+
template <typename T, access::decorated IsDecorated>
39+
void innerFunc(id<1> wiID, global_ptr<T, IsDecorated> ptr_1,
40+
global_ptr<T, IsDecorated> ptr_2,
41+
global_ptr<T, IsDecorated> ptr_3,
42+
global_ptr<T, IsDecorated> ptr_4,
43+
global_ptr<T, IsDecorated> ptr_5,
44+
local_ptr<T, IsDecorated> local_ptr,
45+
private_ptr<T, IsDecorated> priv_ptr) {
4446
T t = ptr_1[wiID.get(0)];
47+
48+
// Write to ptr_2 using local_ptr subscript.
4549
local_ptr[wiID.get(0)] = t;
46-
t = local_ptr[wiID.get(0)];
47-
ptr_2[wiID.get(0)] = t;
50+
ptr_2[wiID.get(0)] = local_ptr[wiID.get(0)];
51+
52+
// Reset local ptr
53+
local_ptr[wiID.get(0)] = 0;
54+
55+
// Write to ptr_3 using dereferencing.
56+
*(local_ptr + wiID.get(0)) = t;
57+
*(ptr_3 + wiID.get(0)) = *(local_ptr + wiID.get(0));
58+
59+
// Write to ptr_2 using priv_ptr subscript.
60+
priv_ptr[wiID.get(0)] = t;
61+
ptr_4[wiID.get(0)] = priv_ptr[wiID.get(0)];
62+
63+
// Reset local ptr
64+
priv_ptr[wiID.get(0)] = 0;
65+
66+
// Write to ptr_3 using dereferencing.
67+
*(priv_ptr + wiID.get(0)) = t;
68+
*(ptr_5 + wiID.get(0)) = *(priv_ptr + wiID.get(0));
4869
}
4970

50-
template <typename T> void testMultPtr() {
71+
template <typename T, access::decorated IsDecorated> void testMultPtr() {
5172
T data_1[10];
52-
for (size_t i = 0; i < 10; ++i) {
53-
data_1[i] = 1;
54-
}
5573
T data_2[10];
74+
T data_3[10];
75+
T data_4[10];
76+
T data_5[10];
5677
for (size_t i = 0; i < 10; ++i) {
78+
data_1[i] = 1;
5779
data_2[i] = 2;
80+
data_3[i] = 3;
81+
data_4[i] = 4;
82+
data_5[i] = 5;
5883
}
5984

6085
{
6186
range<1> numOfItems{10};
6287
buffer<T, 1> bufferData_1(data_1, numOfItems);
6388
buffer<T, 1> bufferData_2(data_2, numOfItems);
89+
buffer<T, 1> bufferData_3(data_3, numOfItems);
90+
buffer<T, 1> bufferData_4(data_4, numOfItems);
91+
buffer<T, 1> bufferData_5(data_5, numOfItems);
6492
queue myQueue;
6593
myQueue.submit([&](handler &cgh) {
6694
accessor<T, 1, access::mode::read, access::target::device,
@@ -69,113 +97,159 @@ template <typename T> void testMultPtr() {
6997
accessor<T, 1, access::mode::read_write, access::target::device,
7098
access::placeholder::false_t>
7199
accessorData_2(bufferData_2, cgh);
100+
accessor<T, 1, access::mode::read_write, access::target::device,
101+
access::placeholder::false_t>
102+
accessorData_3(bufferData_3, cgh);
103+
accessor<T, 1, access::mode::read_write, access::target::device,
104+
access::placeholder::false_t>
105+
accessorData_4(bufferData_4, cgh);
106+
accessor<T, 1, access::mode::read_write, access::target::device,
107+
access::placeholder::false_t>
108+
accessorData_5(bufferData_5, cgh);
72109
local_accessor<T, 1> localAccessor(numOfItems, cgh);
73110

74-
cgh.parallel_for<class testMultPtrKernel<T>>(range<1>{10}, [=](id<1> wiID) {
75-
auto ptr_1 = make_ptr<T, access::address_space::global_space>(
76-
accessorData_1.get_pointer());
77-
auto ptr_2 = make_ptr<T, access::address_space::global_space>(
78-
accessorData_2.get_pointer());
79-
auto local_ptr = make_ptr<T, access::address_space::local_space>(
80-
localAccessor.get_pointer());
81-
82-
// General conversions in multi_ptr class
83-
T *RawPtr = nullptr;
84-
global_ptr<T> ptr_4(RawPtr);
85-
ptr_4 = RawPtr;
86-
87-
global_ptr<T> ptr_5(accessorData_1);
88-
89-
global_ptr<void> ptr_6((void *)RawPtr);
90-
91-
ptr_6 = (void *)RawPtr;
92-
93-
// Explicit conversions for device_ptr/host_ptr to global_ptr
94-
device_ptr<void> ptr_7((void *)RawPtr);
95-
global_ptr<void> ptr_8 = global_ptr<void>(ptr_7);
96-
host_ptr<void> ptr_9((void *)RawPtr);
97-
global_ptr<void> ptr_10 = global_ptr<void>(ptr_9);
98-
// TODO: need propagation of a7b763b26 patch to acl tool before testing
99-
// these conversions - otherwise the test would fail on accelerator
100-
// device during reversed translation from SPIR-V to LLVM IR
101-
// device_ptr<T> ptr_11(accessorData_1);
102-
// global_ptr<T> ptr_12 = global_ptr<T>(ptr_11);
103-
104-
innerFunc<T>(wiID.get(0), ptr_1, ptr_2, local_ptr);
105-
});
111+
cgh.parallel_for<class testMultPtrKernel<T, IsDecorated>>(
112+
range<1>{10}, [=](id<1> wiID) {
113+
T private_data[10];
114+
for (size_t i = 0; i < 10; ++i)
115+
private_data[i] = 0;
116+
localAccessor[wiID] = 0;
117+
118+
auto ptr_1 =
119+
multi_ptr<T, access::address_space::global_space, IsDecorated>(
120+
accessorData_1);
121+
auto ptr_2 =
122+
multi_ptr<T, access::address_space::global_space, IsDecorated>(
123+
accessorData_2);
124+
auto ptr_3 =
125+
multi_ptr<T, access::address_space::global_space, IsDecorated>(
126+
accessorData_3);
127+
auto ptr_4 =
128+
multi_ptr<T, access::address_space::global_space, IsDecorated>(
129+
accessorData_4);
130+
auto ptr_5 =
131+
multi_ptr<T, access::address_space::global_space, IsDecorated>(
132+
accessorData_5);
133+
auto local_ptr =
134+
multi_ptr<T, access::address_space::local_space, IsDecorated>(
135+
localAccessor);
136+
auto priv_ptr =
137+
address_space_cast<access::address_space::private_space,
138+
IsDecorated>(private_data);
139+
static_assert(
140+
std::is_same_v<private_ptr<T, IsDecorated>, decltype(priv_ptr)>,
141+
"Incorrect type for priv_ptr.");
142+
143+
// General conversions in multi_ptr class
144+
T *RawPtr = nullptr;
145+
global_ptr<T, IsDecorated> ptr_6 =
146+
address_space_cast<access::address_space::global_space,
147+
IsDecorated>(RawPtr);
148+
149+
global_ptr<T, IsDecorated> ptr_7(accessorData_1);
150+
151+
global_ptr<void, IsDecorated> ptr_8 =
152+
address_space_cast<access::address_space::global_space,
153+
IsDecorated>((void *)RawPtr);
154+
155+
// Explicit conversions for device_ptr/host_ptr to global_ptr
156+
device_ptr<void, IsDecorated> ptr_9 = address_space_cast<
157+
access::address_space::ext_intel_global_device_space,
158+
IsDecorated>((void *)RawPtr);
159+
global_ptr<void, IsDecorated> ptr_10 =
160+
global_ptr<void, IsDecorated>(ptr_9);
161+
host_ptr<void, IsDecorated> ptr_11 = address_space_cast<
162+
access::address_space::ext_intel_global_host_space,
163+
IsDecorated>((void *)RawPtr);
164+
global_ptr<void, IsDecorated> ptr_12 =
165+
global_ptr<void, IsDecorated>(ptr_11);
166+
167+
innerFunc<T, IsDecorated>(wiID.get(0), ptr_1, ptr_2, ptr_3, ptr_4,
168+
ptr_5, local_ptr, priv_ptr);
169+
});
106170
});
107171
}
108172
for (size_t i = 0; i < 10; ++i) {
109173
assert(data_1[i] == 1 && "Expected data_1[i] == 1");
110-
}
111-
for (size_t i = 0; i < 10; ++i) {
112174
assert(data_2[i] == 1 && "Expected data_2[i] == 1");
175+
assert(data_3[i] == 1 && "Expected data_3[i] == 1");
176+
assert(data_4[i] == 1 && "Expected data_4[i] == 1");
177+
assert(data_5[i] == 1 && "Expected data_5[i] == 1");
113178
}
114179
}
115180

116-
template <typename T> void testMultPtrArrowOperator() {
181+
template <typename T, access::decorated IsDecorated>
182+
void testMultPtrArrowOperator() {
117183
point<T> data_1[1] = {1};
118184
point<T> data_2[1] = {2};
119185
point<T> data_3[1] = {3};
120-
point<T> data_4[1] = {4};
121186

122187
{
123188
range<1> numOfItems{1};
124189
buffer<point<T>, 1> bufferData_1(data_1, numOfItems);
125190
buffer<point<T>, 1> bufferData_2(data_2, numOfItems);
126191
buffer<point<T>, 1> bufferData_3(data_3, numOfItems);
127-
buffer<point<T>, 1> bufferData_4(data_4, numOfItems);
128192
queue myQueue;
129193
myQueue.submit([&](handler &cgh) {
130194
accessor<point<T>, 1, access::mode::read, access::target::device,
131195
access::placeholder::false_t>
132196
accessorData_1(bufferData_1, cgh);
133-
accessor<point<T>, 1, access::mode::read, access::target::constant_buffer,
134-
access::placeholder::false_t>
135-
accessorData_2(bufferData_2, cgh);
136-
local_accessor<point<T>, 1> accessorData_3(1, cgh);
197+
local_accessor<point<T>, 1> accessorData_2(1, cgh);
137198
accessor<point<T>, 1, access::mode::read, access::target::device,
138199
access::placeholder::false_t>
139-
accessorData_4(bufferData_4, cgh);
140-
141-
cgh.single_task<class testMultPtrArrowOperatorKernel<T>>([=]() {
142-
auto ptr_1 = make_ptr<point<T>, access::address_space::global_space>(
143-
accessorData_1.get_pointer());
144-
auto ptr_2 = make_ptr<point<T>, access::address_space::constant_space>(
145-
accessorData_2.get_pointer());
146-
auto ptr_3 = make_ptr<point<T>, access::address_space::local_space>(
147-
accessorData_3.get_pointer());
148-
auto ptr_4 =
149-
make_ptr<point<T>,
150-
access::address_space::ext_intel_global_device_space>(
151-
accessorData_4.get_pointer());
152-
153-
auto x1 = ptr_1 -> x;
154-
auto x2 = ptr_2 -> x;
155-
auto x3 = ptr_3 -> x;
156-
auto x4 = ptr_4 -> x;
157-
158-
static_assert(std::is_same<decltype(x1), T>::value,
159-
"Expected decltype(ptr_1->x) == T");
160-
static_assert(std::is_same<decltype(x2), T>::value,
161-
"Expected decltype(ptr_2->x) == T");
162-
static_assert(std::is_same<decltype(x3), T>::value,
163-
"Expected decltype(ptr_3->x) == T");
164-
static_assert(std::is_same<decltype(x4), T>::value,
165-
"Expected decltype(ptr_4->x) == T");
166-
});
200+
accessorData_3(bufferData_3, cgh);
201+
202+
cgh.single_task<class testMultPtrArrowOperatorKernel<T, IsDecorated>>(
203+
[=]() {
204+
point<T> private_val = 0;
205+
;
206+
207+
auto ptr_1 =
208+
multi_ptr<point<T>, access::address_space::global_space,
209+
IsDecorated>(accessorData_1);
210+
auto ptr_2 = multi_ptr<point<T>, access::address_space::local_space,
211+
IsDecorated>(accessorData_2);
212+
auto ptr_3 =
213+
multi_ptr<point<T>,
214+
access::address_space::ext_intel_global_device_space,
215+
IsDecorated>(accessorData_3);
216+
auto ptr_4 =
217+
address_space_cast<access::address_space::private_space,
218+
IsDecorated>(&private_val);
219+
static_assert(std::is_same_v<private_ptr<point<T>, IsDecorated>,
220+
decltype(ptr_4)>,
221+
"Incorrect type for ptr_4.");
222+
223+
auto x1 = ptr_1->x;
224+
auto x2 = ptr_2->x;
225+
auto x3 = ptr_3->x;
226+
auto x4 = ptr_4->x;
227+
228+
static_assert(std::is_same<decltype(x1), T>::value,
229+
"Expected decltype(ptr_1->x) == T");
230+
static_assert(std::is_same<decltype(x2), T>::value,
231+
"Expected decltype(ptr_2->x) == T");
232+
static_assert(std::is_same<decltype(x3), T>::value,
233+
"Expected decltype(ptr_3->x) == T");
234+
static_assert(std::is_same<decltype(x4), T>::value,
235+
"Expected decltype(ptr_4->x) == T");
236+
});
167237
});
168238
}
169239
}
170240

171-
int main() {
172-
testMultPtr<int>();
173-
testMultPtr<float>();
174-
testMultPtr<point<int>>();
175-
testMultPtr<point<float>>();
241+
template <access::decorated IsDecorated> void runTestsForDecoration() {
242+
testMultPtr<int, IsDecorated>();
243+
testMultPtr<float, IsDecorated>();
244+
testMultPtr<point<int>, IsDecorated>();
245+
testMultPtr<point<float>, IsDecorated>();
176246

177-
testMultPtrArrowOperator<int>();
178-
testMultPtrArrowOperator<float>();
247+
testMultPtrArrowOperator<int, IsDecorated>();
248+
testMultPtrArrowOperator<float, IsDecorated>();
249+
}
179250

251+
int main() {
252+
runTestsForDecoration<access::decorated::yes>();
253+
runTestsForDecoration<access::decorated::no>();
180254
return 0;
181255
}

0 commit comments

Comments
 (0)