Skip to content

Commit 8185f30

Browse files
lbushi25lbushi25
authored
Rewrite tests that fail when usm_shared_allocations not supported. (#12636)
Nearly 40 E2E tests use malloc_shared and do not expect an exception to be thrown in case of failure which is not what the spec says, namely, malloc_shared should throw when the usm_shared_allocations aspect is not supported by the device. At the moment malloc_shared simply returns nullptr, however, once the implementation of malloc_shared is changed to throw an exception these tests will produce failures. This PR is the first of a series of PR's to rewrite these tests so that they don't fail once malloc_shared is changed to conform to the spec. These changes will only affect devices that do not support the relevant aspect. This is done either: - by having llvm-lit require that the aspect is supported for tests that rely heavily on malloc_shared, - by disabling, through if statements, sections of code that use malloc_shared for tests that use it sparsely, - by completely rewriting the malloc_shared code with equivalent buffer/accessor code that is guaranteed to work on all platforms. The latter approach is taken on tests that do not have the purpose of testing malloc_shared but simply use it to allocate memory for the test. These changes aim to eliminate test failures while maintaining test coverage. Subsequent PR's will omit this description and refer to this PR instead.
1 parent 5f1d98a commit 8185f30

11 files changed

+185
-194
lines changed

‎sycl/test-e2e/Annotated_arg_ptr/annotated_arg.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
// RUN: %{build} -o %t.out
22
// RUN: %{run} %t.out
3-
//
3+
// REQUIRES: aspect-usm_shared_allocations
44

55
#include "common.hpp"
66

‎sycl/test-e2e/Annotated_arg_ptr/annotated_ptr.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
// RUN: %{build} -o %t.out
22
// RUN: %{run} %t.out
3-
//
3+
// REQUIRES: aspect-usm_shared_allocations
44

55
#include "common.hpp"
66

‎sycl/test-e2e/Annotated_usm/annotated_usm_kind.cpp

Lines changed: 18 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -124,26 +124,27 @@ template <typename T> void testUsmKind(sycl::queue &q) {
124124
[&]() { return TAnnotated(dev, Ctx, properties{usm_kind_host}); },
125125
[&]() { return ATHost(1, q); }, [&]() { return ATHost(1, Ctx); },
126126
[&]() { return ATAnnotated(1, dev, Ctx, alloc::host); }});
127-
128-
CheckUsmKindAll(
129-
alloc::shared,
130-
std::tuple{
131-
[&]() { return MShared(q); }, [&]() { return MShared(dev, Ctx); },
132-
[&]() { return MAnnotated(dev, Ctx, alloc::shared); },
133-
[&]() { return MAnnotated(dev, Ctx, properties{usm_kind_shared}); },
134-
[&]() { return AShared(1, q); },
135-
[&]() { return AShared(1, dev, Ctx); },
136-
[&]() { return AAnnotated(1, dev, Ctx, alloc::shared); },
137-
[&]() { return TShared(q); }, [&]() { return TShared(dev, Ctx); },
138-
[&]() { return TAnnotated(dev, Ctx, alloc::shared); },
139-
[&]() { return TAnnotated(dev, Ctx, properties{usm_kind_shared}); },
140-
[&]() { return ATShared(1, q); },
141-
[&]() { return ATShared(1, dev, Ctx); },
142-
[&]() { return ATAnnotated(1, dev, Ctx, alloc::shared); }});
127+
if (dev.has(sycl::aspect::usm_shared_allocations)) {
128+
CheckUsmKindAll(
129+
alloc::shared,
130+
std::tuple{
131+
[&]() { return MShared(q); }, [&]() { return MShared(dev, Ctx); },
132+
[&]() { return MAnnotated(dev, Ctx, alloc::shared); },
133+
[&]() { return MAnnotated(dev, Ctx, properties{usm_kind_shared}); },
134+
[&]() { return AShared(1, q); },
135+
[&]() { return AShared(1, dev, Ctx); },
136+
[&]() { return AAnnotated(1, dev, Ctx, alloc::shared); },
137+
[&]() { return TShared(q); }, [&]() { return TShared(dev, Ctx); },
138+
[&]() { return TAnnotated(dev, Ctx, alloc::shared); },
139+
[&]() { return TAnnotated(dev, Ctx, properties{usm_kind_shared}); },
140+
[&]() { return ATShared(1, q); },
141+
[&]() { return ATShared(1, dev, Ctx); },
142+
[&]() { return ATAnnotated(1, dev, Ctx, alloc::shared); }});
143+
}
143144
}
144145

145146
int main() {
146147
sycl::queue q;
147148
testUsmKind<int>(q);
148149
return 0;
149-
}
150+
}

‎sycl/test-e2e/Basic/group_local_memory.cpp

Lines changed: 15 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -7,17 +7,21 @@ constexpr int N = 5;
77

88
int main() {
99
sycl::queue q;
10-
int *ptr = sycl::malloc_shared<int>(N, q);
11-
q.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1> it) {
12-
auto g = it.get_group();
13-
auto mem = sycl::ext::oneapi::group_local_memory<int[N]>(g, 1, 2, 3, 4, 5);
14-
auto ref = *mem;
15-
for (int i = 0; i < N; ++i) {
16-
ptr[i] = ref[i];
17-
}
18-
}).wait();
10+
sycl::buffer<int> buf{sycl::range{N}};
11+
q.submit([&](sycl::handler &h) {
12+
sycl::accessor acc{buf, h};
13+
h.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1> it) {
14+
auto g = it.get_group();
15+
auto mem =
16+
sycl::ext::oneapi::group_local_memory<int[N]>(g, 1, 2, 3, 4, 5);
17+
auto ref = *mem;
18+
for (int i = 0; i < N; ++i) {
19+
acc[i] = ref[i];
20+
}
21+
});
22+
});
23+
sycl::host_accessor result{buf};
1924
for (int i = 0; i < N; ++i) {
20-
assert(ptr[i] == (i + 1));
25+
assert(result[i] == (i + 1));
2126
}
22-
sycl::free(ptr, q);
2327
}

‎sycl/test-e2e/Basic/large-range.cpp

Lines changed: 25 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -26,34 +26,38 @@ void check_sum(std::string_view desc, const ContainerT &data, size_t N) {
2626
template <typename RangeT>
2727
void test_regular(std::string_view desc, queue &q, size_t B, RangeT range) {
2828
auto N = range.size();
29-
std::vector accumulators_v(B, 0, usm_allocator<int, usm::alloc::shared>(q));
30-
auto *accumulators = accumulators_v.data();
31-
32-
q.parallel_for(range, [=](auto it) {
33-
atomic_ref<int, memory_order::relaxed, memory_scope::device> ref(
34-
accumulators[it.get_linear_id() % B]);
35-
++ref;
36-
}).wait();
37-
29+
std::vector accumulators_v(B, 0);
30+
{
31+
sycl::buffer accumulator_buf{accumulators_v};
32+
q.submit([&](sycl::handler &h) {
33+
sycl::accessor accumulators{accumulator_buf, h};
34+
h.parallel_for(range, [=](auto it) {
35+
atomic_ref<int, memory_order::relaxed, memory_scope::device> ref(
36+
accumulators[it.get_linear_id() % B]);
37+
++ref;
38+
});
39+
});
40+
} // destruction of accumulator_buf here writes back data to accumulators_v
3841
check_sum(desc, accumulators_v, N);
3942
}
4043

4144
template <typename RangeT>
4245
void test_spec_constant(std::string_view desc, queue &q, size_t B,
4346
RangeT range) {
4447
auto N = range.size();
45-
std::vector accumulators_v(B, 0, usm_allocator<int, usm::alloc::shared>(q));
46-
auto *accumulators = accumulators_v.data();
47-
48-
q.submit([&](handler &cgh) {
49-
cgh.set_specialization_constant<C>(2);
50-
cgh.parallel_for(range, [=](auto it, kernel_handler h) {
51-
atomic_ref<int, memory_order::relaxed, memory_scope::device> ref(
52-
accumulators[it.get_linear_id() % B]);
53-
ref += h.get_specialization_constant<C>();
54-
});
55-
}).wait();
56-
48+
std::vector accumulators_v(B, 0);
49+
{
50+
sycl::buffer accumulators_buf{accumulators_v};
51+
q.submit([&](handler &cgh) {
52+
sycl::accessor accumulators{accumulators_buf, cgh};
53+
cgh.set_specialization_constant<C>(2);
54+
cgh.parallel_for(range, [=](auto it, kernel_handler h) {
55+
atomic_ref<int, memory_order::relaxed, memory_scope::device> ref(
56+
accumulators[it.get_linear_id() % B]);
57+
ref += h.get_specialization_constant<C>();
58+
});
59+
});
60+
} // destruction of accumulators_buf here writes data back to accumulators_v
5761
check_sum(desc, accumulators_v, N * 2);
5862
}
5963

‎sycl/test-e2e/Basic/span.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
//
44
// Fails to release USM pointer on HIP for NVIDIA
55
// XFAIL: hip_nvidia
6-
6+
// REQUIRES: aspect-usm_shared_allocations
77
#include <numeric>
88
#include <sycl/sycl.hpp>
99

‎sycl/test-e2e/Basic/wrapped_usm_pointers.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
// RUN: %{build} -o %t.out
22
// RUN: %{run} %t.out
3-
3+
// REQUIRES: aspect-usm_shared_allocations
44
//==---------- wrapped_usm_pointer.cpp - test pointers in struct ---------==//
55
//
66
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

‎sycl/test-e2e/Complex/sycl_complex_math_test.cpp

Lines changed: 50 additions & 62 deletions
Original file line numberDiff line numberDiff line change
@@ -13,31 +13,26 @@
1313
bool operator()(sycl::queue &Q, cmplx<T> init, \
1414
cmplx<T> ref = cmplx<T>(0, 0), bool use_ref = false) { \
1515
bool pass = true; \
16-
\
1716
auto std_in = init_std_complex(init.re, init.im); \
1817
experimental::complex<T> cplx_input{init.re, init.im}; \
19-
\
20-
auto *cplx_out = sycl::malloc_shared<experimental::complex<T>>(1, Q); \
21-
\
18+
sycl::buffer<experimental::complex<T>> cplx_out_buf{sycl::range{1}}; \
2219
/*Get std::complex output*/ \
2320
std::complex<T> std_out{ref.re, ref.im}; \
2421
if (!use_ref) \
2522
std_out = std::math_func(std_in); \
26-
\
2723
/*Check cplx::complex output from device*/ \
28-
Q.single_task([=]() { \
29-
cplx_out[0] = experimental::math_func<T>(cplx_input); \
30-
}).wait(); \
31-
\
32-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ true); \
24+
Q.submit([&](sycl::handler &h) { \
25+
sycl::accessor cplx_out{cplx_out_buf, h}; \
26+
h.single_task( \
27+
[=]() { cplx_out[0] = experimental::math_func<T>(cplx_input); }); \
28+
}); \
29+
sycl::host_accessor cplx_out_acc{cplx_out_buf}; \
30+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ true); \
3331
\
3432
/*Check cplx::complex output from host*/ \
35-
cplx_out[0] = experimental::math_func<T>(cplx_input); \
36-
\
37-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ false); \
38-
\
39-
sycl::free(cplx_out, Q); \
33+
cplx_out_acc[0] = experimental::math_func<T>(cplx_input); \
4034
\
35+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ false); \
4136
return pass; \
4237
} \
4338
};
@@ -73,28 +68,26 @@ TEST_MATH_OP_TYPE(tanh)
7368
\
7469
auto std_in = init_std_complex(init.re, init.im); \
7570
experimental::complex<T> cplx_input{init.re, init.im}; \
76-
\
77-
auto *cplx_out = sycl::malloc_shared<T>(1, Q); \
71+
sycl::buffer<T> cplx_out_buf{sycl::range{1}}; \
7872
\
7973
/*Get std::complex output*/ \
8074
T std_out = ref.re; \
8175
if (!use_ref) \
8276
std_out = std::math_func(std_in); \
8377
\
8478
/*Check cplx::complex output from device*/ \
85-
Q.single_task([=]() { \
86-
cplx_out[0] = experimental::math_func<T>(cplx_input); \
87-
}).wait(); \
88-
\
89-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ true); \
79+
Q.submit([&](sycl::handler &h) { \
80+
sycl::accessor cplx_out{cplx_out_buf, h}; \
81+
h.single_task( \
82+
[=]() { cplx_out[0] = experimental::math_func<T>(cplx_input); }); \
83+
}); \
84+
sycl::host_accessor cplx_out_acc{cplx_out_buf}; \
85+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ true); \
9086
\
9187
/*Check cplx::complex output from host*/ \
92-
cplx_out[0] = experimental::math_func<T>(cplx_input); \
93-
\
94-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ false); \
95-
\
96-
sycl::free(cplx_out, Q); \
88+
cplx_out_acc[0] = experimental::math_func<T>(cplx_input); \
9789
\
90+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ false); \
9891
return pass; \
9992
} \
10093
};
@@ -121,23 +114,21 @@ TEST_MATH_OP_TYPE(imag)
121114
std::complex<T> std_out = ref; \
122115
if (!use_ref) \
123116
std_out = std::math_func(std_in); \
124-
\
125-
auto *cplx_out = sycl::malloc_shared<experimental::complex<T>>(1, Q); \
126-
\
117+
sycl::buffer<experimental::complex<T>> cplx_out_buf{sycl::range{1}}; \
127118
/*Check cplx::complex output from device*/ \
128-
Q.single_task([=]() { \
129-
cplx_out[0] = experimental::math_func<X>(std_in); \
130-
}).wait(); \
119+
Q.submit([&](sycl::handler &h) { \
120+
sycl::accessor cplx_out{cplx_out_buf, h}; \
121+
h.single_task( \
122+
[=]() { cplx_out[0] = experimental::math_func<X>(std_in); }); \
123+
}); \
124+
sycl::host_accessor cplx_out_acc{cplx_out_buf}; \
131125
\
132-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ true); \
126+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ true); \
133127
\
134128
/*Check cplx::complex output from host*/ \
135-
cplx_out[0] = experimental::math_func<X>(std_in); \
136-
\
137-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ false); \
138-
\
139-
sycl::free(cplx_out, Q); \
129+
cplx_out_acc[0] = experimental::math_func<X>(std_in); \
140130
\
131+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ false); \
141132
return pass; \
142133
} \
143134
};
@@ -161,23 +152,21 @@ TEST_MATH_OP_TYPE(proj)
161152
T std_out = ref; \
162153
if (!use_ref) \
163154
std_out = std::math_func(std_in); \
164-
\
165-
auto *cplx_out = sycl::malloc_shared<T>(1, Q); \
166-
\
155+
sycl::buffer<T> cplx_out_buf{sycl::range{1}}; \
167156
/*Check cplx::complex output from device*/ \
168-
Q.single_task([=]() { \
169-
cplx_out[0] = experimental::math_func<X>(init); \
170-
}).wait(); \
157+
Q.submit([&](sycl::handler &h) { \
158+
sycl::accessor cplx_out{cplx_out_buf, h}; \
159+
h.single_task( \
160+
[=]() { cplx_out[0] = experimental::math_func<X>(std_in); }); \
161+
}); \
162+
sycl::host_accessor cplx_out_acc{cplx_out_buf}; \
171163
\
172-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ true); \
164+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ true); \
173165
\
174166
/*Check cplx::complex output from host*/ \
175-
cplx_out[0] = experimental::math_func<X>(init); \
176-
\
177-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ false); \
178-
\
179-
sycl::free(cplx_out, Q); \
167+
cplx_out_acc[0] = experimental::math_func<X>(init); \
180168
\
169+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ false); \
181170
return pass; \
182171
} \
183172
};
@@ -197,26 +186,25 @@ template <typename T> struct test_polar {
197186
bool use_ref = false) {
198187
bool pass = true;
199188

200-
auto *cplx_out = sycl::malloc_shared<experimental::complex<T>>(1, Q);
201-
189+
sycl::buffer<experimental::complex<T>> cplx_out_buf{sycl::range(1)};
202190
/*Get std::complex output*/
203191
std::complex<T> std_out{ref.re, ref.im};
204192
if (!use_ref)
205193
std_out = std::polar(init.re, init.im);
206194

207195
/*Check cplx::complex output from device*/
208-
Q.single_task([=]() {
209-
cplx_out[0] = experimental::polar<T>(init.re, init.im);
210-
}).wait();
211-
212-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ true);
196+
Q.submit([&](sycl::handler &h) {
197+
sycl::accessor cplx_out{cplx_out_buf, h};
198+
h.single_task(
199+
[=]() { cplx_out[0] = experimental::polar<T>(init.re, init.im); });
200+
});
201+
sycl::host_accessor cplx_out_acc{cplx_out_buf};
202+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ true);
213203

214204
/*Check cplx::complex output from host*/
215-
cplx_out[0] = experimental::polar<T>(init.re, init.im);
216-
217-
pass &= check_results(cplx_out[0], std_out, /*is_device*/ false);
205+
cplx_out_acc[0] = experimental::polar<T>(init.re, init.im);
218206

219-
sycl::free(cplx_out, Q);
207+
pass &= check_results(cplx_out_acc[0], std_out, /*is_device*/ false);
220208

221209
return pass;
222210
}

0 commit comments

Comments
 (0)