Skip to content

Commit d591f23

Browse files
vasilytricbb-sycl
authored andcommitted
[SYCL][ESIMD] Add test for simd constructor from vector (intel#687)
1 parent 9755cdd commit d591f23

File tree

1 file changed

+199
-0
lines changed

1 file changed

+199
-0
lines changed
Lines changed: 199 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,199 @@
1+
//==------- ctor_vector_core.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+
// TODO The tests freezed during runtime when using simd<char, 32>. The
18+
// SIMD_RUN_TEST_WITH_CHAR_TYPES macros must be enabled when it is resolved.
19+
//
20+
// TODO This test disabled due to simd<short, 32> vector filled with unexpected
21+
// values from 16th element. The issue was created
22+
// https://github.com/intel/llvm/issues/5245 and and the
23+
// SIMD_RUN_TEST_WITH_VECTOR_LEN_32 macros must be enabled when it is resolved.
24+
//
25+
// Test for simd constructor from vector.
26+
// This test uses different data types, dimensionality and different simd
27+
// constructor invocation contexts.
28+
// The test do the following actions:
29+
// - call init_simd.data() to retreive vector_type and then provide it to the
30+
// simd constructor
31+
// - bitwise comparing expected and retrieved values
32+
33+
#include "common.hpp"
34+
35+
using namespace esimd_test::api::functional;
36+
using namespace sycl::ext::intel::experimental::esimd;
37+
38+
// Descriptor class for the case of calling constructor in initializer context.
39+
struct initializer {
40+
static std::string get_description() { return "initializer"; }
41+
42+
template <typename DataT, int NumElems>
43+
static void call_simd_ctor(const DataT *const ref_data, DataT *const out) {
44+
simd<DataT, NumElems> init_simd(ref_data);
45+
const auto simd_by_init = simd<DataT, NumElems>(init_simd.data());
46+
simd_by_init.copy_to(out);
47+
}
48+
};
49+
50+
// Descriptor class for the case of calling constructor in variable declaration
51+
// context.
52+
struct var_decl {
53+
static std::string get_description() { return "variable declaration"; }
54+
55+
template <typename DataT, int NumElems>
56+
static void call_simd_ctor(const DataT *const ref_data, DataT *const out) {
57+
simd<DataT, NumElems> init_simd(ref_data);
58+
simd<DataT, NumElems> simd_by_var_decl(init_simd.data());
59+
simd_by_var_decl.copy_to(out);
60+
}
61+
};
62+
63+
// Descriptor class for the case of calling constructor in rvalue in an
64+
// expression context.
65+
struct rval_in_expr {
66+
static std::string get_description() { return "rvalue in an expression"; }
67+
68+
template <typename DataT, int NumElems>
69+
static void call_simd_ctor(const DataT *const ref_data, DataT *const out) {
70+
simd<DataT, NumElems> init_simd(ref_data);
71+
simd<DataT, NumElems> simd_by_rval;
72+
simd_by_rval = simd<DataT, NumElems>(init_simd.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(const DataT *const ref_data, DataT *const out) {
85+
simd<DataT, NumElems> init_simd(ref_data);
86+
call_simd_by_const_ref<DataT, NumElems>(
87+
simd<DataT, NumElems>(init_simd.data()), out);
88+
}
89+
90+
private:
91+
template <typename DataT, int NumElems>
92+
static void
93+
call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref,
94+
DataT *out) {
95+
simd_by_const_ref.copy_to(out);
96+
}
97+
};
98+
99+
// The main test routine.
100+
// 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 {
102+
public:
103+
bool operator()(sycl::queue &queue, const std::string &data_type) {
104+
bool passed = true;
105+
const std::vector<DataT> ref_data = generate_ref_data<DataT, NumElems>();
106+
107+
// If current number of elements is equal to one, then run test with each
108+
// one value from reference data.
109+
// If current number of elements is greater than one, then run tests with
110+
// whole reference data.
111+
if constexpr (NumElems == 1) {
112+
for (size_t i = 0; i < ref_data.size(); ++i) {
113+
passed &= run_verification(queue, {ref_data[i]}, data_type);
114+
}
115+
} else {
116+
passed &= run_verification(queue, ref_data, data_type);
117+
}
118+
return passed;
119+
}
120+
121+
private:
122+
bool run_verification(sycl::queue &queue, const std::vector<DataT> &ref_data,
123+
const std::string &data_type) {
124+
assert(ref_data.size() == NumElems &&
125+
"Reference data size is not equal to the simd vector length.");
126+
127+
bool passed = true;
128+
129+
shared_allocator<DataT> allocator(queue);
130+
shared_vector<DataT> result(NumElems, allocator);
131+
shared_vector<DataT> shared_ref_data(ref_data.begin(), ref_data.end(),
132+
allocator);
133+
134+
queue.submit([&](sycl::handler &cgh) {
135+
const DataT *const ref = shared_ref_data.data();
136+
DataT *const out = result.data();
137+
138+
cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
139+
[=]() SYCL_ESIMD_KERNEL {
140+
TestCaseT::template call_simd_ctor<DataT, NumElems>(ref, out);
141+
});
142+
});
143+
queue.wait_and_throw();
144+
145+
for (size_t i = 0; i < result.size(); ++i) {
146+
if (!are_bitwise_equal(ref_data[i], result[i])) {
147+
passed = false;
148+
149+
const auto description =
150+
ctors::TestDescription<DataT, NumElems, TestCaseT>(
151+
i, result[i], ref_data[i], data_type);
152+
log::fail(description);
153+
}
154+
}
155+
156+
return passed;
157+
}
158+
};
159+
160+
int main(int, char **) {
161+
sycl::queue queue(esimd_test::ESIMDSelector{},
162+
esimd_test::createExceptionHandler());
163+
164+
sycl::device device = queue.get_device();
165+
// verify aspect::fp16 due to using sycl::half data type
166+
// verify aspect::fp64 due to using double data type
167+
if (!device.is_host() && !device.has(sycl::aspect::fp16) &&
168+
!device.has(sycl::aspect::fp64)) {
169+
std::cout << "Test skipped\n";
170+
return 0;
171+
}
172+
173+
bool passed = true;
174+
175+
#ifdef SIMD_RUN_TEST_WITH_CHAR_TYPES
176+
const auto types = get_tested_types<tested_types::all>();
177+
#else
178+
const auto types = named_type_pack<short, unsigned short, int, unsigned int,
179+
long, unsigned long, float, double,
180+
long long, unsigned long long>(
181+
{"short", "unsigned short", "int", "unsigned int", "long",
182+
"unsigned long", "float", "double", "long long", "unsigned long long"});
183+
#endif
184+
#ifdef SIMD_RUN_TEST_WITH_VECTOR_LEN_32
185+
const auto dims = get_all_dimensions();
186+
#else
187+
const auto dims = values_pack<1, 8, 16>();
188+
#endif
189+
190+
// Run for specific combinations of types, vector length and invocation
191+
// contexts.
192+
passed &= for_all_types_and_dims<run_test, initializer>(types, dims, queue);
193+
passed &= for_all_types_and_dims<run_test, var_decl>(types, dims, queue);
194+
passed &= for_all_types_and_dims<run_test, rval_in_expr>(types, dims, queue);
195+
passed &= for_all_types_and_dims<run_test, const_ref>(types, dims, queue);
196+
197+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
198+
return passed ? 0 : 1;
199+
}

0 commit comments

Comments
 (0)