Skip to content

Commit 58b0949

Browse files
yuriykochbb-sycl
authored andcommitted
[SYCL][ESIMD] Provide the for_all_combinations utility (intel#721)
[SYCL][ESIMD] Provide the for_all_combinations utility Provides a generic way to represent a combinations of types and compile-time values to cover. Supports any number of type and value packs to iterate over Adds unnamed type pack support By doing so we avoid additional restriction on architecture of the test case functors and on logging architecture, as we don't need to construct string description for every type/value on the top level. We are still free to use named type packs, giving us a possibility to have any combination of named packs and unnamed packs, with value packs supported in both variations. Tests are updated as there is no more for_all_types_and_dims function Signed-off-by: Kochetkov, Yuriy <[email protected]>
1 parent 25255a7 commit 58b0949

File tree

9 files changed

+632
-7
lines changed

9 files changed

+632
-7
lines changed

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

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,149 @@
2525
#include "ctor_array.hpp"
2626

2727
using namespace esimd_test::api::functional;
28+
<<<<<<< HEAD
29+
=======
30+
using namespace sycl::ext::intel::experimental::esimd;
31+
32+
// Descriptor class for the case of calling constructor in initializer context.
33+
struct initializer {
34+
static std::string get_description() { return "initializer"; }
35+
36+
template <typename DataT, int NumElems>
37+
static void call_simd_ctor(DataT (&&ref_data)[NumElems], DataT *const out) {
38+
static_assert(
39+
type_traits::is_nonconst_rvalue_reference_v<decltype(ref_data)>,
40+
"Provided input data is not nonconst rvalue reference");
41+
const auto simd_by_init = simd<DataT, NumElems>(std::move(ref_data));
42+
simd_by_init.copy_to(out);
43+
}
44+
};
45+
46+
// Descriptor class for the case of calling constructor in variable declaration
47+
// context.
48+
struct var_decl {
49+
static std::string get_description() { return "variable declaration"; }
50+
51+
template <typename DataT, int NumElems>
52+
static void call_simd_ctor(DataT (&&ref_data)[NumElems], DataT *const out) {
53+
static_assert(
54+
type_traits::is_nonconst_rvalue_reference_v<decltype(ref_data)>,
55+
"Provided input data is not nonconst rvalue reference");
56+
const simd<DataT, NumElems> simd_by_var_decl(std::move(ref_data));
57+
simd_by_var_decl.copy_to(out);
58+
}
59+
};
60+
61+
// Descriptor class for the case of calling constructor in rvalue in an
62+
// expression context.
63+
struct rval_in_expr {
64+
static std::string get_description() { return "rvalue in an expression"; }
65+
66+
template <typename DataT, int NumElems>
67+
static void call_simd_ctor(DataT (&&ref_data)[NumElems], DataT *const out) {
68+
static_assert(
69+
type_traits::is_nonconst_rvalue_reference_v<decltype(ref_data)>,
70+
"Provided input data is not nonconst rvalue reference");
71+
simd<DataT, NumElems> simd_by_rval;
72+
simd_by_rval = simd<DataT, NumElems>(std::move(ref_data));
73+
simd_by_rval.copy_to(out);
74+
}
75+
};
76+
77+
// Descriptor class for the case of calling constructor in const reference
78+
// context.
79+
class const_ref {
80+
public:
81+
static std::string get_description() { return "const reference"; }
82+
83+
template <typename DataT, int NumElems>
84+
static void call_simd_ctor(DataT (&&ref_data)[NumElems], DataT *const out) {
85+
static_assert(
86+
type_traits::is_nonconst_rvalue_reference_v<decltype(ref_data)>,
87+
"Provided input data is not nonconst rvalue reference");
88+
call_simd_by_const_ref<DataT, NumElems>(
89+
simd<DataT, NumElems>(std::move(ref_data)), out);
90+
}
91+
92+
private:
93+
template <typename DataT, int NumElems>
94+
static void
95+
call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref,
96+
DataT *out) {
97+
simd_by_const_ref.copy_to(out);
98+
}
99+
};
100+
101+
// The main test routine.
102+
// Using functor class to be able to iterate over the pre-defined data types.
103+
template <typename DataT, typename DimT, typename TestCaseT> class run_test {
104+
static constexpr int NumElems = DimT::value;
105+
106+
public:
107+
bool operator()(sycl::queue &queue, const std::string &data_type) {
108+
109+
bool passed = true;
110+
const std::vector<DataT> ref_data = generate_ref_data<DataT, NumElems>();
111+
112+
// If current number of elements is equal to one, then run test with each
113+
// one value from reference data.
114+
// If current number of elements is greater than one, then run tests with
115+
// whole reference data.
116+
if constexpr (NumElems == 1) {
117+
for (size_t i = 0; i < ref_data.size(); ++i) {
118+
passed &= run_verification(queue, {ref_data[i]}, data_type);
119+
}
120+
} else {
121+
passed &= run_verification(queue, ref_data, data_type);
122+
}
123+
return passed;
124+
}
125+
126+
private:
127+
bool run_verification(sycl::queue &queue, const std::vector<DataT> &ref_data,
128+
const std::string &data_type) {
129+
assert(ref_data.size() == NumElems &&
130+
"Reference data size is not equal to the simd vector length.");
131+
132+
bool passed = true;
133+
134+
shared_allocator<DataT> allocator(queue);
135+
shared_vector<DataT> result(NumElems, allocator);
136+
shared_vector<DataT> shared_ref_data(ref_data.begin(), ref_data.end(),
137+
allocator);
138+
139+
queue.submit([&](sycl::handler &cgh) {
140+
const DataT *const ref = shared_ref_data.data();
141+
DataT *const out = result.data();
142+
143+
cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
144+
[=]() SYCL_ESIMD_KERNEL {
145+
DataT ref_on_dev[NumElems];
146+
for (size_t i = 0; i < NumElems; ++i) {
147+
ref_on_dev[i] = ref[i];
148+
}
149+
150+
TestCaseT::template call_simd_ctor<DataT, NumElems>(
151+
std::move(ref_on_dev), out);
152+
});
153+
});
154+
queue.wait_and_throw();
155+
156+
for (size_t i = 0; i < result.size(); ++i) {
157+
if (!are_bitwise_equal(ref_data[i], result[i])) {
158+
passed = false;
159+
160+
const auto description =
161+
ctors::TestDescription<DataT, NumElems, TestCaseT>(
162+
i, result[i], ref_data[i], data_type);
163+
log::fail(description);
164+
}
165+
}
166+
167+
return passed;
168+
}
169+
};
170+
>>>>>>> 6870ea3ee ([SYCL][ESIMD] Provide the for_all_combinations utility (#721))
28171

29172
int main(int, char **) {
30173
sycl::queue queue(esimd_test::ESIMDSelector{},
@@ -41,10 +184,17 @@ int main(int, char **) {
41184
=======
42185
const auto types = get_tested_types<tested_types::all>();
43186
const auto dims = get_all_dimensions();
187+
<<<<<<< HEAD
44188
>>>>>>> 7ff842b8f ([SYCL][ESIMD] Enable verification for 32 simd vector length (#758))
45189

46190
passed &=
47191
for_all_combinations<ctors::run_test>(types, sizes, contexts, queue);
192+
=======
193+
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
194+
const_ref>::generate();
195+
196+
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
197+
>>>>>>> 6870ea3ee ([SYCL][ESIMD] Provide the for_all_combinations utility (#721))
48198

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

SYCL/ESIMD/api/functional/ctors/ctor_copy.hpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,9 +102,14 @@ class const_ref {
102102

103103
// The main test routine.
104104
// Using functor class to be able to iterate over the pre-defined data types.
105+
<<<<<<< HEAD:SYCL/ESIMD/api/functional/ctors/ctor_copy.hpp
105106
template <typename DataT, typename SizeT, typename TestCaseT> class run_test {
106107
static constexpr int NumElems = SizeT::value;
107108
using TestDescriptionT = ctors::TestDescription<NumElems, TestCaseT>;
109+
=======
110+
template <typename DataT, typename DimT, typename TestCaseT> class run_test {
111+
static constexpr int NumElems = DimT::value;
112+
>>>>>>> 6870ea3ee ([SYCL][ESIMD] Provide the for_all_combinations utility (#721)):SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
108113

109114
public:
110115
bool operator()(sycl::queue &queue, const std::string &data_type) {
@@ -170,4 +175,23 @@ template <typename DataT, typename SizeT, typename TestCaseT> class run_test {
170175
}
171176
};
172177

178+
<<<<<<< HEAD:SYCL/ESIMD/api/functional/ctors/ctor_copy.hpp
173179
} // namespace esimd_test::api::functional::ctors
180+
=======
181+
int main(int, char **) {
182+
sycl::queue queue(esimd_test::ESIMDSelector{},
183+
esimd_test::createExceptionHandler());
184+
185+
bool passed = true;
186+
187+
const auto types = get_tested_types<tested_types::all>();
188+
const auto dims = get_all_dimensions();
189+
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
190+
const_ref>::generate();
191+
192+
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
193+
194+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
195+
return passed ? 0 : 1;
196+
}
197+
>>>>>>> 6870ea3ee ([SYCL][ESIMD] Provide the for_all_combinations utility (#721)):SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
Lines changed: 131 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,131 @@
1+
//==------- ctor_default.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+
// REQUIRES: gpu, level_zero
9+
// XREQUIRES: gpu
10+
// TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
11+
// The current "REQUIRES" should be replaced with "gpu" only as mentioned in
12+
// "XREQUIRES".
13+
// UNSUPPORTED: cuda, hip
14+
// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
15+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
16+
//
17+
// Test for esimd default constructor.
18+
//
19+
// The simd can't be constructed with sycl::half data type. The issue was
20+
// created (https://github.com/intel/llvm/issues/5077) and the TEST_HALF macros
21+
// must be enabled when it is resolved.
22+
23+
#include "common.hpp"
24+
25+
using namespace sycl::ext::intel::experimental::esimd;
26+
using namespace esimd_test::api::functional;
27+
28+
// Descriptor class for the case of calling constructor in initializer context
29+
struct initializer {
30+
static std::string get_description() { return "initializer"; }
31+
32+
template <typename DataT, int NumElems>
33+
static void call_simd_ctor(DataT *const output_data) {
34+
const auto simd_by_init = simd<DataT, NumElems>();
35+
simd_by_init.copy_to(output_data);
36+
}
37+
};
38+
39+
// Descriptor class for the case of calling constructor in variable declaration
40+
// context
41+
struct var_decl {
42+
static std::string get_description() { return "variable declaration"; }
43+
44+
template <typename DataT, int NumElems>
45+
static void call_simd_ctor(DataT *const output_data) {
46+
simd<DataT, NumElems> simd_by_var_decl;
47+
simd_by_var_decl.copy_to(output_data);
48+
}
49+
};
50+
51+
// Descriptor class for the case of calling constructor in rvalue in an
52+
// expression context
53+
struct rval_in_expr {
54+
static std::string get_description() { return "rvalue in an expression"; }
55+
56+
template <typename DataT, int NumElems>
57+
static void call_simd_ctor(DataT *const output_data) {
58+
simd<DataT, NumElems> simd_by_rval;
59+
simd_by_rval = simd<DataT, NumElems>();
60+
simd_by_rval.copy_to(output_data);
61+
}
62+
};
63+
64+
// Descriptor class for the case of calling constructor in const reference
65+
// context
66+
struct const_ref {
67+
static std::string get_description() { return "const reference"; }
68+
69+
template <typename DataT, int NumElems>
70+
static void
71+
call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref,
72+
DataT *output_data) {
73+
simd_by_const_ref.copy_to(output_data);
74+
}
75+
76+
template <typename DataT, int NumElems>
77+
static void call_simd_ctor(DataT *const output_data) {
78+
call_simd_by_const_ref<DataT, NumElems>(simd<DataT, NumElems>(),
79+
output_data);
80+
}
81+
};
82+
83+
// Struct that calls simd in provided context and then verifies obtained result.
84+
template <typename DataT, typename DimT, typename TestCaseT> struct run_test {
85+
static constexpr int NumElems = DimT::value;
86+
87+
bool operator()(sycl::queue &queue, const std::string &data_type) {
88+
bool passed = true;
89+
DataT default_val{};
90+
91+
shared_vector<DataT> result(NumElems, shared_allocator<DataT>(queue));
92+
93+
queue.submit([&](sycl::handler &cgh) {
94+
DataT *const out = result.data();
95+
cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
96+
[=]() SYCL_ESIMD_KERNEL {
97+
TestCaseT::template call_simd_ctor<DataT, NumElems>(out);
98+
});
99+
});
100+
101+
for (size_t i = 0; i < result.size(); ++i) {
102+
if (result[i] != default_val) {
103+
passed = false;
104+
105+
const auto description =
106+
ctors::TestDescription<DataT, NumElems, TestCaseT>(
107+
i, result[i], default_val, data_type);
108+
log::fail(description);
109+
}
110+
}
111+
112+
return passed;
113+
}
114+
};
115+
116+
int main(int, char **) {
117+
sycl::queue queue(esimd_test::ESIMDSelector{},
118+
esimd_test::createExceptionHandler());
119+
120+
bool passed = true;
121+
122+
const auto types = get_tested_types<tested_types::all>();
123+
const auto dims = get_all_dimensions();
124+
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
125+
const_ref>::generate();
126+
127+
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
128+
129+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
130+
return passed ? 0 : 1;
131+
}

0 commit comments

Comments
 (0)