Skip to content

Commit c2d204a

Browse files
vasilytricbb-sycl
vasilytric
authored and
bb-sycl
committed
[SYCL][ESIMD] Add tests on simd increment and decrement operators (intel#827)
* [SYCL][ESIMD] Add tests on simd increment and decrement operators
1 parent bfbb9c0 commit c2d204a

5 files changed

+493
-0
lines changed

‎SYCL/ESIMD/api/functional/operators/operator_decrement_and_increment.hpp

Lines changed: 296 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,296 @@
1+
//===-- operator_decrement_and_increment.hpp - Functions for tests on simd
2+
// increment and decrement operators definition. ---------------------===//
3+
//
4+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5+
// See https://llvm.org/LICENSE.txt for license information.
6+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7+
//
8+
//===----------------------------------------------------------------------===//
9+
///
10+
/// \file
11+
/// This file provides functions for tests on simd decrement and increment
12+
/// operators.
13+
///
14+
//===----------------------------------------------------------------------===//
15+
16+
#pragma once
17+
18+
#include "../mutator.hpp"
19+
#include "common.hpp"
20+
21+
namespace esimd = sycl::ext::intel::experimental::esimd;
22+
namespace esimd_functional = esimd_test::api::functional;
23+
24+
namespace esimd_test::api::functional::operators {
25+
26+
// Descriptor class for the case of calling constructor in initializer context.
27+
struct pre_decrement {
28+
static std::string get_description() { return "pre decrement"; }
29+
30+
template <typename DataT, int NumElems>
31+
static void call_simd_ctor(const DataT *const ref_data,
32+
DataT *const source_simd_out,
33+
DataT *const result_simd_out) {
34+
auto source_simd = esimd::simd<DataT, NumElems>();
35+
source_simd.copy_from(ref_data);
36+
esimd::simd<DataT, NumElems> result_simd = --source_simd;
37+
source_simd.copy_to(source_simd_out);
38+
result_simd.copy_to(result_simd_out);
39+
}
40+
41+
template <typename DataT> static DataT apply_operator(DataT &val) {
42+
return --val;
43+
}
44+
45+
static constexpr bool is_increment() { return false; }
46+
};
47+
48+
// Descriptor class for the case of calling constructor in initializer context.
49+
struct post_decrement {
50+
static std::string get_description() { return "post decrement"; }
51+
52+
template <typename DataT, int NumElems>
53+
static void call_simd_ctor(const DataT *const ref_data,
54+
DataT *const source_simd_out,
55+
DataT *const result_simd_out) {
56+
auto source_simd = esimd::simd<DataT, NumElems>();
57+
source_simd.copy_from(ref_data);
58+
esimd::simd<DataT, NumElems> result_simd = source_simd--;
59+
source_simd.copy_to(source_simd_out);
60+
result_simd.copy_to(result_simd_out);
61+
}
62+
63+
template <typename DataT> static DataT apply_operator(DataT &val) {
64+
return val--;
65+
}
66+
67+
static constexpr bool is_increment() { return false; }
68+
};
69+
70+
// Descriptor class for the case of calling constructor in initializer context.
71+
struct pre_increment {
72+
static std::string get_description() { return "pre increment"; }
73+
74+
template <typename DataT, int NumElems>
75+
static void call_simd_ctor(const DataT *const ref_data,
76+
DataT *const source_simd_out,
77+
DataT *const result_simd_out) {
78+
auto source_simd = esimd::simd<DataT, NumElems>();
79+
source_simd.copy_from(ref_data);
80+
esimd::simd<DataT, NumElems> result_simd = ++source_simd;
81+
source_simd.copy_to(source_simd_out);
82+
result_simd.copy_to(result_simd_out);
83+
}
84+
85+
template <typename DataT> static DataT apply_operator(DataT &val) {
86+
return ++val;
87+
}
88+
89+
static constexpr bool is_increment() { return true; }
90+
};
91+
92+
// Descriptor class for the case of calling constructor in initializer context.
93+
struct post_increment {
94+
static std::string get_description() { return "post increment"; }
95+
96+
template <typename DataT, int NumElems>
97+
static void call_simd_ctor(const DataT *const ref_data,
98+
DataT *const source_simd_out,
99+
DataT *const result_simd_out) {
100+
auto source_simd = esimd::simd<DataT, NumElems>();
101+
source_simd.copy_from(ref_data);
102+
esimd::simd<DataT, NumElems> result_simd = source_simd++;
103+
source_simd.copy_to(source_simd_out);
104+
result_simd.copy_to(result_simd_out);
105+
}
106+
107+
template <typename DataT> static DataT apply_operator(DataT &val) {
108+
return val++;
109+
}
110+
111+
static constexpr bool is_increment() { return true; }
112+
};
113+
114+
template <typename DataT, int NumElems, typename TestCaseT>
115+
class IncrementAndDecrementTestDescription : public ITestDescription {
116+
public:
117+
IncrementAndDecrementTestDescription(size_t index, DataT retrieved_val,
118+
DataT expected_val,
119+
const std::string &error_details,
120+
const std::string &data_type)
121+
: m_data_type(data_type), m_retrieved_val(retrieved_val),
122+
m_expected_val(expected_val), m_index(index),
123+
m_error_details(error_details) {}
124+
125+
std::string to_string() const override {
126+
std::string log_msg("Failed for simd<");
127+
128+
log_msg += m_data_type + ", " + std::to_string(NumElems) + ">";
129+
log_msg += ", retrieved: " + std::to_string(m_retrieved_val);
130+
log_msg += ", expected: " + std::to_string(m_expected_val);
131+
log_msg += ", at index: " + std::to_string(m_index);
132+
log_msg += " for " + TestCaseT::get_description() + " operator: ";
133+
log_msg += m_error_details;
134+
135+
return log_msg;
136+
}
137+
138+
private:
139+
const std::string m_data_type;
140+
const DataT m_retrieved_val;
141+
const DataT m_expected_val;
142+
const size_t m_index;
143+
const std::string m_error_details;
144+
};
145+
146+
struct base_test {
147+
template <typename DataT, int NumElems, typename TestCaseT>
148+
static std::vector<DataT> generate_input_data() {
149+
std::vector<DataT> ref_data = generate_ref_data<DataT, NumElems>();
150+
151+
if constexpr (TestCaseT::is_increment()) {
152+
mutate(ref_data, mutator::For_addition<DataT>(1));
153+
} else {
154+
mutate(ref_data, mutator::For_subtraction<DataT>(1));
155+
}
156+
157+
return ref_data;
158+
}
159+
};
160+
161+
struct fp_accuracy_test {
162+
template <typename DataT, int NumElems, typename TestCaseT>
163+
static std::vector<DataT> generate_input_data() {
164+
std::vector<DataT> ref_data;
165+
166+
static const DataT min = value<DataT>::lowest();
167+
static const DataT denorm_min = value<DataT>::denorm_min();
168+
static const DataT max = value<DataT>::max();
169+
static const DataT inexact = static_cast<DataT>(0.1);
170+
171+
if constexpr (TestCaseT::is_increment()) {
172+
ref_data.reserve((NumElems > 1) ? NumElems : 6);
173+
ref_data.insert(ref_data.end(),
174+
{inexact, denorm_min, -denorm_min,
175+
value<DataT>::pos_ulp(static_cast<DataT>(-1.0)),
176+
value<DataT>::pos_ulp(min),
177+
value<DataT>::neg_ulp(max - 1)});
178+
179+
} else {
180+
ref_data.reserve((NumElems > 1) ? NumElems : 6);
181+
ref_data.insert(ref_data.end(),
182+
{inexact, denorm_min, -denorm_min,
183+
value<DataT>::neg_ulp(static_cast<DataT>(-1.0)),
184+
value<DataT>::neg_ulp(max),
185+
value<DataT>::pos_ulp(min + 1)});
186+
}
187+
188+
for (size_t i = ref_data.size(); i < NumElems; ++i) {
189+
ref_data.push_back(inexact * i);
190+
}
191+
192+
return ref_data;
193+
}
194+
};
195+
196+
// The main test routine.
197+
// Using functor class to be able to iterate over the pre-defined data types.
198+
template <typename IsAccuracyTestT, typename DataT, typename SizeT,
199+
typename TestCaseT>
200+
class run_test {
201+
static constexpr int NumElems = SizeT::value;
202+
203+
public:
204+
bool operator()(sycl::queue &queue, const std::string &data_type) {
205+
bool passed = true;
206+
std::vector<DataT> ref_data =
207+
IsAccuracyTestT::template generate_input_data<DataT, NumElems,
208+
TestCaseT>();
209+
210+
// If current number of elements is equal to one, then run test with each
211+
// one value from reference data.
212+
// If current number of elements is greater than one, then run tests with
213+
// whole reference data.
214+
if constexpr (NumElems == 1) {
215+
for (size_t i = 0; i < ref_data.size(); ++i) {
216+
passed &= run_verification(queue, {ref_data[i]}, data_type);
217+
}
218+
} else {
219+
passed &= run_verification(queue, ref_data, data_type);
220+
}
221+
return passed;
222+
}
223+
224+
private:
225+
bool run_verification(sycl::queue &queue, const std::vector<DataT> &ref_data,
226+
const std::string &data_type) {
227+
assert(ref_data.size() == NumElems &&
228+
"Reference data size is not equal to the simd vector length.");
229+
230+
bool passed = true;
231+
232+
shared_allocator<DataT> allocator(queue);
233+
shared_vector<DataT> source_simd_out(NumElems, allocator);
234+
shared_vector<DataT> result_simd_out(NumElems, allocator);
235+
shared_vector<DataT> shared_ref_data(ref_data.begin(), ref_data.end(),
236+
allocator);
237+
238+
queue.submit([&](sycl::handler &cgh) {
239+
const DataT *const ref = shared_ref_data.data();
240+
DataT *source_simd_data_ptr = source_simd_out.data();
241+
DataT *result_simd_data_ptr = result_simd_out.data();
242+
243+
cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
244+
[=]() SYCL_ESIMD_KERNEL {
245+
TestCaseT::template call_simd_ctor<DataT, NumElems>(
246+
ref, source_simd_data_ptr, result_simd_data_ptr);
247+
});
248+
});
249+
queue.wait_and_throw();
250+
251+
for (size_t i = 0; i < NumElems; ++i) {
252+
DataT expected_source_value = shared_ref_data[i];
253+
254+
const DataT expected_return_value =
255+
TestCaseT::apply_operator(expected_source_value);
256+
257+
passed &= verify_result(i, expected_source_value, source_simd_out[i],
258+
"unexpected argument modification", data_type);
259+
passed &= verify_result(i, expected_return_value, result_simd_out[i],
260+
"unexpected return value", data_type);
261+
}
262+
263+
return passed;
264+
}
265+
266+
bool verify_result(size_t i, DataT expected, DataT retrieved,
267+
const std::string &simd_type,
268+
const std::string &data_type) {
269+
bool passed = true;
270+
if constexpr (type_traits::is_sycl_floating_point_v<DataT>) {
271+
if (std::isnan(expected) && !std::isnan(retrieved)) {
272+
passed = false;
273+
274+
// TODO: Make ITestDescription architecture more flexible.
275+
// We are assuming that the NaN opcode may differ
276+
std::string log_msg("Failed for simd<");
277+
log_msg += data_type + ", " + std::to_string(NumElems) + ">";
278+
log_msg += ". The element at index: " + std::to_string(i) +
279+
", is not nan, but it should.";
280+
281+
log::note(log_msg);
282+
}
283+
}
284+
if (!are_bitwise_equal(expected, retrieved)) {
285+
passed = false;
286+
287+
const auto description =
288+
IncrementAndDecrementTestDescription<DataT, NumElems, TestCaseT>(
289+
i, retrieved, expected, simd_type, data_type);
290+
log::fail(description);
291+
}
292+
return passed;
293+
}
294+
};
295+
296+
} // namespace esimd_test::api::functional::operators

‎SYCL/ESIMD/api/functional/operators/operator_decrement_and_increment_accuracy_core.cpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
//==------- operator_decrement_and_increment_accuracy_core.cpp - DPC++ ESIMD
2+
// on-device test -------------------------------------------------==//
3+
//
4+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5+
// See https://llvm.org/LICENSE.txt for license information.
6+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7+
//
8+
//===----------------------------------------------------------------------===//
9+
// REQUIRES: gpu, level_zero
10+
// XREQUIRES: gpu
11+
// TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
12+
// The current "REQUIRES" should be replaced with "gpu" only as mentioned in
13+
// "XREQUIRES".
14+
// UNSUPPORTED: cuda, hip
15+
// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
16+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
17+
//
18+
// Test for simd increment and decrement operators.
19+
// The test creates source simd instance and call increment or decrement
20+
// operator opera with reference data. The test verifies that in the output data
21+
// contained correctness data according to chosen operator and has no precision
22+
// differences with interaction with floating point data types.
23+
24+
#include "operator_decrement_and_increment.hpp"
25+
26+
using namespace esimd_test::api::functional;
27+
28+
int main(int, char **) {
29+
sycl::queue queue(esimd_test::ESIMDSelector{},
30+
esimd_test::createExceptionHandler());
31+
32+
bool passed = true;
33+
34+
const auto fp_types = get_tested_types<tested_types::fp>();
35+
const auto all_sizes = get_all_sizes();
36+
const auto contexts =
37+
unnamed_type_pack<operators::pre_increment,
38+
operators::post_increment>::generate();
39+
40+
passed &=
41+
for_all_combinations<operators::run_test, operators::fp_accuracy_test>(
42+
fp_types, all_sizes, contexts, queue);
43+
44+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
45+
return passed ? 0 : 1;
46+
}

‎SYCL/ESIMD/api/functional/operators/operator_decrement_and_increment_accuracy_fp_extra.cpp

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
//==------- operator_decrement_and_increment_accuracy_fp_extra.cpp - DPC++
2+
// ESIMD on-device test -------------------------------------------==//
3+
//
4+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5+
// See https://llvm.org/LICENSE.txt for license information.
6+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7+
//
8+
//===----------------------------------------------------------------------===//
9+
// REQUIRES: gpu, level_zero
10+
// XREQUIRES: gpu
11+
// TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
12+
// The current "REQUIRES" should be replaced with "gpu" only as mentioned in
13+
// "XREQUIRES".
14+
// UNSUPPORTED: cuda, hip
15+
// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
16+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
17+
//
18+
// TODO simd<sycl::half, N> vector filled with unexpected values. The
19+
// ESIMD_TESTS_RUN_WITH_HALF macros must be enabled when it is resolved.
20+
//
21+
// Test for simd increment and decrement operators.
22+
// The test creates source simd instance and call increment or decrement
23+
// operator opera with reference data. The test verifies that in the output data
24+
// contained correctness data according to chosen operator and has no precision
25+
// differences with interaction with floating point data types.
26+
27+
#include "operator_decrement_and_increment.hpp"
28+
29+
using namespace esimd_test::api::functional;
30+
31+
int main(int, char **) {
32+
sycl::queue queue(esimd_test::ESIMDSelector{},
33+
esimd_test::createExceptionHandler());
34+
35+
bool passed = true;
36+
37+
#ifdef ESIMD_TESTS_RUN_WITH_HALF
38+
const auto fp_extra_types = get_tested_types<tested_types::fp_extra>();
39+
#else
40+
const auto fp_extra_types = named_type_pack<double>::generate("double");
41+
#endif
42+
const auto all_sizes = get_all_sizes();
43+
const auto contexts =
44+
unnamed_type_pack<operators::pre_increment, operators::post_increment,
45+
operators::pre_decrement,
46+
operators::post_decrement>::generate();
47+
48+
passed &=
49+
for_all_combinations<operators::run_test, operators::fp_accuracy_test>(
50+
fp_extra_types, all_sizes, contexts, queue);
51+
52+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
53+
return passed ? 0 : 1;
54+
}

0 commit comments

Comments
 (0)