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

Commit a4c51cc

Browse files
dongkyunahn-inteldongkyunahn-intel
committed
Merge branch 'intel' of https://github.com/intel/llvm-test-suite into invoke_esimd_emulator_check_in
2 parents bbdb607 + af11f41 commit a4c51cc

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

57 files changed

+723
-566
lines changed

‎SYCL/Basic/buffer/buffer_create.cpp

100755100644
Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,11 +11,15 @@ int main() {
1111
constexpr int Size = 100;
1212
queue Queue;
1313
auto D = Queue.get_device();
14-
14+
auto NumOfDevices = Queue.get_context().get_devices().size();
1515
buffer<::cl_int, 1> Buffer(Size);
1616
Queue.submit([&](handler &cgh) {
1717
accessor Accessor{Buffer, cgh, read_write};
18-
if (D.get_info<info::device::host_unified_memory>())
18+
if (NumOfDevices > 1)
19+
// Currently the Level Zero plugin uses host allocations for multi-device
20+
// contexts because such allocations are accessible by all devices.
21+
std::cerr << "Multi GPU should use zeMemAllocHost\n";
22+
else if (D.get_info<info::device::host_unified_memory>())
1923
std::cerr << "Integrated GPU should use zeMemAllocHost\n";
2024
else
2125
std::cerr << "Discrete GPU should use zeMemAllocDevice\n";
@@ -26,5 +30,5 @@ int main() {
2630
return 0;
2731
}
2832

29-
// CHECK: {{Integrated|Discrete}} GPU should use [[API:zeMemAllocHost|zeMemAllocDevice]]
33+
// CHECK: {{Integrated|Multi|Discrete}} GPU should use [[API:zeMemAllocHost|zeMemAllocHost|zeMemAllocDevice]]
3034
// CHECK: ZE ---> [[API]](

‎SYCL/Basic/image/image_sample.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,8 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
22
// RUN: %CPU_RUN_PLACEHOLDER %t.out
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
4-
// UNSUPPORTED: cuda || hip
4+
// Temporarily disable test on Windows due to regressions in GPU driver.
5+
// UNSUPPORTED: cuda || hip, windows
56

67
#include <CL/sycl.hpp>
78

‎SYCL/ESIMD/PrefixSum.cpp

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -199,21 +199,21 @@ void cmk_acum_final(unsigned *buf, unsigned h_pos, unsigned int stride_elems,
199199
cnt_table.select<1, 1, 8, 4>(j, 3) += cnt_table.select<1, 1, 8, 4>(j, 1);
200200
// step 3
201201
cnt_table.select<1, 1, 4, 1>(j, 4) +=
202-
cnt_table.replicate<1, 0, 4, 0>(j, 3);
202+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 3);
203203
cnt_table.select<1, 1, 4, 1>(j, 12) +=
204-
cnt_table.replicate<1, 0, 4, 0>(j, 11);
204+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 11);
205205
cnt_table.select<1, 1, 4, 1>(j, 20) +=
206-
cnt_table.replicate<1, 0, 4, 0>(j, 19);
206+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 19);
207207
cnt_table.select<1, 1, 4, 1>(j, 28) +=
208-
cnt_table.replicate<1, 0, 4, 0>(j, 27);
208+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 27);
209209
// step 4
210210
cnt_table.select<1, 1, 8, 1>(j, 8) +=
211-
cnt_table.replicate<1, 0, 8, 0>(j, 7);
211+
cnt_table.replicate_vs_w_hs<1, 0, 8, 0>(j, 7);
212212
cnt_table.select<1, 1, 8, 1>(j, 24) +=
213-
cnt_table.replicate<1, 0, 8, 0>(j, 23);
213+
cnt_table.replicate_vs_w_hs<1, 0, 8, 0>(j, 23);
214214
// step 5
215215
cnt_table.select<1, 1, 16, 1>(j, 16) +=
216-
cnt_table.replicate<1, 0, 16, 0>(j, 15);
216+
cnt_table.replicate_vs_w_hs<1, 0, 16, 0>(j, 15);
217217
}
218218
scatter_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset, S,
219219
p);
@@ -268,21 +268,21 @@ void cmk_prefix_iterative(unsigned *buf, unsigned h_pos,
268268
cnt_table.select<1, 1, 8, 4>(j, 3) += cnt_table.select<1, 1, 8, 4>(j, 1);
269269
// step 3
270270
cnt_table.select<1, 1, 4, 1>(j, 4) +=
271-
cnt_table.replicate<1, 0, 4, 0>(j, 3);
271+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 3);
272272
cnt_table.select<1, 1, 4, 1>(j, 12) +=
273-
cnt_table.replicate<1, 0, 4, 0>(j, 11);
273+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 11);
274274
cnt_table.select<1, 1, 4, 1>(j, 20) +=
275-
cnt_table.replicate<1, 0, 4, 0>(j, 19);
275+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 19);
276276
cnt_table.select<1, 1, 4, 1>(j, 28) +=
277-
cnt_table.replicate<1, 0, 4, 0>(j, 27);
277+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 27);
278278
// step 4
279279
cnt_table.select<1, 1, 8, 1>(j, 8) +=
280-
cnt_table.replicate<1, 0, 8, 0>(j, 7);
280+
cnt_table.replicate_vs_w_hs<1, 0, 8, 0>(j, 7);
281281
cnt_table.select<1, 1, 8, 1>(j, 24) +=
282-
cnt_table.replicate<1, 0, 8, 0>(j, 23);
282+
cnt_table.replicate_vs_w_hs<1, 0, 8, 0>(j, 23);
283283
// step 5
284284
cnt_table.select<1, 1, 16, 1>(j, 16) +=
285-
cnt_table.replicate<1, 0, 16, 0>(j, 15);
285+
cnt_table.replicate_vs_w_hs<1, 0, 16, 0>(j, 15);
286286
}
287287

288288
// during reduction phase, we've already computed prefix sum and saved in

‎SYCL/ESIMD/Prefix_Local_sum3.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -211,21 +211,21 @@ void cmk_acum_final(unsigned *buf, unsigned h_pos, unsigned int stride_elems,
211211
cnt_table.select<1, 1, 8, 4>(j, 3) += cnt_table.select<1, 1, 8, 4>(j, 1);
212212
// step 3
213213
cnt_table.select<1, 1, 4, 1>(j, 4) +=
214-
cnt_table.replicate<1, 0, 4, 0>(j, 3);
214+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 3);
215215
cnt_table.select<1, 1, 4, 1>(j, 12) +=
216-
cnt_table.replicate<1, 0, 4, 0>(j, 11);
216+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 11);
217217
cnt_table.select<1, 1, 4, 1>(j, 20) +=
218-
cnt_table.replicate<1, 0, 4, 0>(j, 19);
218+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 19);
219219
cnt_table.select<1, 1, 4, 1>(j, 28) +=
220-
cnt_table.replicate<1, 0, 4, 0>(j, 27);
220+
cnt_table.replicate_vs_w_hs<1, 0, 4, 0>(j, 27);
221221
// step 4
222222
cnt_table.select<1, 1, 8, 1>(j, 8) +=
223-
cnt_table.replicate<1, 0, 8, 0>(j, 7);
223+
cnt_table.replicate_vs_w_hs<1, 0, 8, 0>(j, 7);
224224
cnt_table.select<1, 1, 8, 1>(j, 24) +=
225-
cnt_table.replicate<1, 0, 8, 0>(j, 23);
225+
cnt_table.replicate_vs_w_hs<1, 0, 8, 0>(j, 23);
226226
// step 5
227227
cnt_table.select<1, 1, 16, 1>(j, 16) +=
228-
cnt_table.replicate<1, 0, 16, 0>(j, 15);
228+
cnt_table.replicate_vs_w_hs<1, 0, 16, 0>(j, 15);
229229
}
230230
scatter_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset, S,
231231
p);

‎SYCL/ESIMD/api/functional/ctors/ctor_array_core.cpp

Lines changed: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -98,7 +98,9 @@ class const_ref {
9898

9999
// The main test routine.
100100
// Using functor class to be able to iterate over the pre-defined data types.
101-
template <typename DataT, int NumElems, typename TestCaseT> class run_test {
101+
template <typename DataT, typename DimT, typename TestCaseT> class run_test {
102+
static constexpr int NumElems = DimT::value;
103+
102104
public:
103105
bool operator()(sycl::queue &queue, const std::string &data_type) {
104106

@@ -172,13 +174,10 @@ int main(int, char **) {
172174

173175
const auto types = get_tested_types<tested_types::all>();
174176
const auto dims = get_all_dimensions();
177+
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
178+
const_ref>::generate();
175179

176-
// Run for specific combinations of types, vector length, and invocation
177-
// contexts.
178-
passed &= for_all_types_and_dims<run_test, initializer>(types, dims, queue);
179-
passed &= for_all_types_and_dims<run_test, var_decl>(types, dims, queue);
180-
passed &= for_all_types_and_dims<run_test, rval_in_expr>(types, dims, queue);
181-
passed &= for_all_types_and_dims<run_test, const_ref>(types, dims, queue);
180+
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
182181

183182
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
184183
return passed ? 0 : 1;

‎SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -88,7 +88,9 @@ class const_ref {
8888

8989
// The main test routine.
9090
// Using functor class to be able to iterate over the pre-defined data types.
91-
template <typename DataT, int NumElems, typename TestCaseT> class run_test {
91+
template <typename DataT, typename DimT, typename TestCaseT> class run_test {
92+
static constexpr int NumElems = DimT::value;
93+
9294
public:
9395
bool operator()(sycl::queue &queue, const std::string &data_type) {
9496
bool passed = true;
@@ -155,11 +157,10 @@ int main(int, char **) {
155157

156158
const auto types = get_tested_types<tested_types::all>();
157159
const auto dims = get_all_dimensions();
160+
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
161+
const_ref>::generate();
158162

159-
passed &= for_all_types_and_dims<run_test, initializer>(types, dims, queue);
160-
passed &= for_all_types_and_dims<run_test, var_decl>(types, dims, queue);
161-
passed &= for_all_types_and_dims<run_test, rval_in_expr>(types, dims, queue);
162-
passed &= for_all_types_and_dims<run_test, const_ref>(types, dims, queue);
163+
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
163164

164165
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
165166
return passed ? 0 : 1;

‎SYCL/ESIMD/api/functional/ctors/ctor_default.cpp

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -81,7 +81,9 @@ struct const_ref {
8181
};
8282

8383
// Struct that calls simd in provided context and then verifies obtained result.
84-
template <typename DataT, int NumElems, typename TestCaseT> struct run_test {
84+
template <typename DataT, typename DimT, typename TestCaseT> struct run_test {
85+
static constexpr int NumElems = DimT::value;
86+
8587
bool operator()(sycl::queue &queue, const std::string &data_type) {
8688
bool passed = true;
8789
DataT default_val{};
@@ -119,11 +121,10 @@ int main(int, char **) {
119121

120122
const auto types = get_tested_types<tested_types::all>();
121123
const auto dims = get_all_dimensions();
124+
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
125+
const_ref>::generate();
122126

123-
passed &= for_all_types_and_dims<run_test, initializer>(types, dims, queue);
124-
passed &= for_all_types_and_dims<run_test, var_decl>(types, dims, queue);
125-
passed &= for_all_types_and_dims<run_test, rval_in_expr>(types, dims, queue);
126-
passed &= for_all_types_and_dims<run_test, const_ref>(types, dims, queue);
127+
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
127128

128129
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
129130
return passed ? 0 : 1;

‎SYCL/ESIMD/api/functional/ctors/ctor_fill.hpp

Lines changed: 24 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@ struct initializer {
3636

3737
// Descriptor class for the case of calling constructor in variable declaration
3838
// context.
39-
struct var_dec {
39+
struct var_decl {
4040
static std::string get_description() { return "variable declaration"; }
4141

4242
template <typename DataT, int NumElems>
@@ -48,7 +48,7 @@ struct var_dec {
4848

4949
// Descriptor class for the case of calling constructor in rvalue in an
5050
// expression context.
51-
struct rval_in_express {
51+
struct rval_in_expr {
5252
static std::string get_description() { return "rvalue in an expression"; }
5353

5454
template <typename DataT, int NumElems>
@@ -191,30 +191,32 @@ class FillCtorTestDescription
191191
}
192192
};
193193

194-
template <typename DataT, int NumElems, typename TestCaseT, typename BaseVal,
195-
typename Step>
194+
template <init_val... Values> auto get_init_values_pack() {
195+
return value_pack<init_val, Values...>::generate_unnamed();
196+
}
197+
198+
template <typename DataT, typename DimT, typename TestCaseT, typename BaseValT,
199+
typename StepT>
196200
class run_test {
201+
static constexpr int NumElems = DimT::value;
202+
static constexpr init_val BaseVal = BaseValT::value;
203+
static constexpr init_val Step = StepT::value;
204+
using KernelT = kernel_for_fill<DataT, NumElems, TestCaseT, BaseVal, Step>;
205+
197206
public:
198207
bool operator()(sycl::queue &queue, const std::string &data_type) {
199-
static_assert(std::is_same_v<typename BaseVal::value_type, init_val>,
200-
"BaseVal template parameter should be init_val type.");
201-
static_assert(std::is_same_v<typename Step::value_type, init_val>,
202-
"Step template parameter should be init_val type.");
203-
204208
shared_vector<DataT> result(NumElems, shared_allocator<DataT>(queue));
205209

206-
const auto base_value = get_value<DataT, BaseVal::value>();
207-
const auto step_value = get_value<DataT, Step::value>(base_value);
210+
const auto base_value = get_value<DataT, BaseVal>();
211+
const auto step_value = get_value<DataT, Step>(base_value);
208212

209213
queue.submit([&](sycl::handler &cgh) {
210214
DataT *const out = result.data();
211215

212-
cgh.single_task<kernel_for_fill<DataT, NumElems, TestCaseT,
213-
BaseVal::value, Step::value>>(
214-
[=]() SYCL_ESIMD_KERNEL {
215-
TestCaseT::template call_simd_ctor<DataT, NumElems>(
216-
base_value, step_value, out);
217-
});
216+
cgh.single_task<KernelT>([=]() SYCL_ESIMD_KERNEL {
217+
TestCaseT::template call_simd_ctor<DataT, NumElems>(base_value,
218+
step_value, out);
219+
});
218220
});
219221
queue.wait_and_throw();
220222
bool passed = true;
@@ -228,8 +230,7 @@ class run_test {
228230
// constructor.
229231
DataT expected_value = base_value;
230232
for (size_t i = 1; i < result.size(); ++i) {
231-
if constexpr (BaseVal::value == init_val::nan ||
232-
Step::value == init_val::nan) {
233+
if constexpr (BaseVal == init_val::nan || Step == init_val::nan) {
233234

234235
if (!std::isnan(result[i])) {
235236
passed = false;
@@ -241,9 +242,8 @@ class run_test {
241242
log_msg += ", with context: " + TestCaseT::get_description();
242243
log_msg += ". The element at index: " + std::to_string(i) +
243244
", is not nan, but it should.";
244-
log_msg +=
245-
", with base value: " + init_val_to_string<BaseVal::value>();
246-
log_msg += ", with step value: " + init_val_to_string<Step::value>();
245+
log_msg += ", with base value: " + init_val_to_string<BaseVal>();
246+
log_msg += ", with step value: " + init_val_to_string<Step>();
247247

248248
log::note(log_msg);
249249
}
@@ -262,33 +262,12 @@ class run_test {
262262
bool fail_test(size_t index, DataT retrieved, DataT expected,
263263
const std::string &data_type) {
264264
const auto description =
265-
FillCtorTestDescription<DataT, NumElems, TestCaseT, BaseVal::value,
266-
Step::value>(index, retrieved, expected,
267-
data_type);
265+
FillCtorTestDescription<DataT, NumElems, TestCaseT, BaseVal, Step>(
266+
index, retrieved, expected, data_type);
268267
log::fail(description);
269268

270269
return false;
271270
}
272271
};
273272

274-
// Iterating over provided types and dimensions, running test for each of
275-
// them.
276-
template <typename TestT, init_val BaseVal, init_val Step, typename... Types,
277-
int... Dims>
278-
bool run_verification(
279-
sycl::queue &queue,
280-
const esimd_functional::values_pack<Dims...> &dimensions,
281-
const esimd_functional::named_type_pack<Types...> &types) {
282-
283-
typedef std::integral_constant<init_val, BaseVal> base_value;
284-
typedef std::integral_constant<init_val, Step> step_value;
285-
286-
bool passed = true;
287-
passed &= esimd_functional::for_all_types_and_dims<run_test, TestT,
288-
base_value, step_value>(
289-
types, dimensions, queue);
290-
291-
return passed;
292-
}
293-
294273
} // namespace esimd_test::api::functional::ctors

‎SYCL/ESIMD/api/functional/ctors/ctor_fill_accuracy_core.cpp

Lines changed: 20 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -28,38 +28,38 @@
2828

2929
using namespace sycl::ext::intel::experimental::esimd;
3030
using namespace esimd_test::api::functional;
31+
using init_val = ctors::init_val;
3132

3233
int main(int, char **) {
3334
sycl::queue queue(esimd_test::ESIMDSelector{},
3435
esimd_test::createExceptionHandler());
3536

3637
bool passed = true;
3738

38-
const auto fp_types = get_tested_types<tested_types::fp>();
39-
const auto single_dim = values_pack<8>();
39+
// Using single dimension and context to verify the accuracy of operations
40+
// with floating point data types
41+
const auto types = get_tested_types<tested_types::fp>();
42+
const auto dims = get_dimensions<8>();
43+
const auto contexts = unnamed_type_pack<ctors::var_decl>::generate();
4044

4145
// Run for specific combinations of types, base and step values and vector
4246
// length.
43-
// The first init_val value it's a base value and the second init_val value
44-
// it's a step value.
4547
#ifdef SIMD_RUN_TEST_WITH_DENORM_INIT_VAL_AND_ULP_STEP
46-
passed &= ctors::run_verification<ctors::var_dec, ctors::init_val::denorm,
47-
ctors::init_val::ulp>(queue, single_dim,
48-
fp_types);
48+
{
49+
const auto base_values = ctors::get_init_values_pack<init_val::denorm>();
50+
const auto step_values = ctors::get_init_values_pack<init_val::ulp>();
51+
passed &= for_all_combinations<ctors::run_test>(
52+
types, dims, contexts, base_values, step_values, queue);
53+
}
4954
#endif
50-
passed &= ctors::run_verification<ctors::var_dec, ctors::init_val::inexact,
51-
ctors::init_val::ulp>(queue, single_dim,
52-
fp_types);
53-
passed &= ctors::run_verification<ctors::var_dec, ctors::init_val::min,
54-
ctors::init_val::ulp>(queue, single_dim,
55-
fp_types);
56-
57-
passed &= ctors::run_verification<ctors::var_dec, ctors::init_val::inexact,
58-
ctors::init_val::ulp_half>(
59-
queue, single_dim, fp_types);
60-
passed &= ctors::run_verification<ctors::var_dec, ctors::init_val::min,
61-
ctors::init_val::ulp_half>(
62-
queue, single_dim, fp_types);
55+
{
56+
const auto base_values =
57+
ctors::get_init_values_pack<init_val::inexact, init_val::min>();
58+
const auto step_values =
59+
ctors::get_init_values_pack<init_val::ulp, init_val::ulp_half>();
60+
passed &= for_all_combinations<ctors::run_test>(
61+
types, dims, contexts, base_values, step_values, queue);
62+
}
6363

6464
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
6565
return passed ? 0 : 1;

0 commit comments

Comments
 (0)