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

Commit 44ef815

Browse files
vasilytricvasilytric
authored
[SYCL][ESIMD] Add test on simd converting constructor (#695)
- Change run_test signature - Add add expected test cases - This update required for case when current value is demoralized because the rule of rounding on the host and device sides is undefined. - Verify that one value is nan, instead of verify thaе both values is nan - Add additional constructor calling contexts - Add coverage for 1 simd size - Enable test for fp_extra types - Use new esimd namespace
1 parent 9992b63 commit 44ef815

File tree

3 files changed

+357
-0
lines changed

3 files changed

+357
-0
lines changed

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

Lines changed: 225 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,225 @@
1+
//===-- ctor_converting.hpp - Functions for tests on simd converting constructor
2+
// 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 converting constructor.
12+
///
13+
//===----------------------------------------------------------------------===//
14+
15+
#pragma once
16+
17+
#include "../value_conv.hpp"
18+
#include "common.hpp"
19+
20+
namespace esimd = sycl::ext::intel::esimd;
21+
22+
namespace esimd_test::api::functional::ctors {
23+
24+
// Descriptor class for the case of calling constructor in initializer context.
25+
struct initializer {
26+
static std::string get_description() { return "initializer"; }
27+
28+
template <typename SrcT, typename DstT, int NumElems>
29+
static void call_simd_ctor(const SrcT *const ref_data, DstT *const out) {
30+
esimd::simd<SrcT, NumElems> input_simd;
31+
input_simd.copy_from(ref_data);
32+
33+
esimd::simd<DstT, NumElems> output_simd = input_simd;
34+
output_simd.copy_to(out);
35+
}
36+
};
37+
38+
// Descriptor class for the case of calling constructor in variable declaration
39+
// context.
40+
struct var_decl {
41+
static std::string get_description() { return "variable declaration"; }
42+
43+
template <typename SrcT, typename DstT, int NumElems>
44+
static void call_simd_ctor(const SrcT *const ref_data, DstT *const out) {
45+
esimd::simd<SrcT, NumElems> input_simd;
46+
input_simd.copy_from(ref_data);
47+
48+
esimd::simd<DstT, NumElems> output_simd(input_simd);
49+
output_simd.copy_to(out);
50+
}
51+
};
52+
53+
// Descriptor class for the case of calling constructor in rvalue in an
54+
// expression context.
55+
struct rval_in_expr {
56+
static std::string get_description() { return "rvalue in an expression"; }
57+
58+
template <typename SrcT, typename DstT, int NumElems>
59+
static void call_simd_ctor(const SrcT *const ref_data, DstT *const out) {
60+
esimd::simd<SrcT, NumElems> input_simd;
61+
input_simd.copy_from(ref_data);
62+
63+
esimd::simd<DstT, NumElems> output_simd;
64+
output_simd = esimd::simd<DstT, NumElems>(input_simd);
65+
output_simd.copy_to(out);
66+
}
67+
};
68+
69+
// Descriptor class for the case of calling constructor in const reference
70+
// context.
71+
class const_ref {
72+
public:
73+
static std::string get_description() { return "const reference"; }
74+
75+
template <typename SrcT, typename DstT, int NumElems>
76+
static void call_simd_ctor(const SrcT *const ref_data, DstT *const out) {
77+
esimd::simd<SrcT, NumElems> input_simd;
78+
input_simd.copy_from(ref_data);
79+
call_simd_by_const_ref<SrcT, DstT, NumElems>(
80+
esimd::simd<SrcT, NumElems>(input_simd), out);
81+
}
82+
83+
private:
84+
template <typename SrcT, typename DstT, int NumElems>
85+
static void
86+
call_simd_by_const_ref(const esimd::simd<SrcT, NumElems> &simd_by_const_ref,
87+
DstT *out) {
88+
esimd::simd<DstT, NumElems> output_simd = simd_by_const_ref;
89+
output_simd.copy_to(out);
90+
}
91+
};
92+
93+
template <typename T, int NumElems, typename ContextT>
94+
class ConvCtorTestDescription : public ITestDescription {
95+
public:
96+
ConvCtorTestDescription(size_t index, T retrieved_val, T expected_val,
97+
const std::string &src_data_type,
98+
const std::string &dst_data_type)
99+
: m_src_data_type(src_data_type), m_dst_data_type(dst_data_type),
100+
m_retrieved_val(retrieved_val), m_expected_val(expected_val),
101+
m_index(index) {}
102+
103+
std::string to_string() const override {
104+
// TODO: Make strings for fp values more short during failure output, may be
105+
// by using hex representation
106+
std::string log_msg("Failed for converting from simd<");
107+
108+
log_msg += m_src_data_type + ", " + std::to_string(NumElems) + ">";
109+
log_msg +=
110+
", to simd<" + m_dst_data_type + ", " + std::to_string(NumElems) + ">";
111+
log_msg += ", with context: " + ContextT::get_description();
112+
log_msg += ", retrieved: " + std::to_string(m_retrieved_val);
113+
log_msg += ", expected: " + std::to_string(m_expected_val);
114+
log_msg += ", at index: " + std::to_string(m_index);
115+
116+
return log_msg;
117+
}
118+
119+
private:
120+
const std::string m_src_data_type;
121+
const std::string m_dst_data_type;
122+
const T m_retrieved_val;
123+
const T m_expected_val;
124+
const size_t m_index;
125+
};
126+
127+
// The main test routine.
128+
// Using functor class to be able to iterate over the pre-defined data types.
129+
template <typename SrcT, typename DimT, typename DstT, typename TestCaseT>
130+
class run_test {
131+
static constexpr int NumElems = DimT::value;
132+
133+
public:
134+
bool operator()(sycl::queue &queue, const std::string &src_data_type,
135+
const std::string &dst_data_type) {
136+
bool passed = true;
137+
const std::vector<SrcT> ref_data =
138+
generate_ref_conv_data<SrcT, DstT, NumElems>();
139+
140+
// If current number of elements is equal to one, then run test with each
141+
// one value from reference data.
142+
// If current number of elements is greater than one, then run tests with
143+
// whole reference data.
144+
if constexpr (NumElems == 1) {
145+
for (size_t i = 0; i < ref_data.size(); ++i) {
146+
passed &= run_verification(queue, {ref_data[i]}, src_data_type,
147+
dst_data_type);
148+
}
149+
} else {
150+
passed &= run_verification(queue, ref_data, src_data_type, dst_data_type);
151+
}
152+
return passed;
153+
}
154+
155+
private:
156+
bool run_verification(sycl::queue &queue, const std::vector<SrcT> &ref_data,
157+
const std::string &src_data_type,
158+
const std::string &dst_data_type) {
159+
assert(ref_data.size() == NumElems &&
160+
"Reference data size is not equal to the simd vector length.");
161+
162+
bool passed = true;
163+
164+
shared_vector<DstT> result(NumElems, shared_allocator<DstT>(queue));
165+
shared_vector<SrcT> shared_ref_data(ref_data.begin(), ref_data.end(),
166+
shared_allocator<SrcT>(queue));
167+
168+
queue.submit([&](sycl::handler &cgh) {
169+
const SrcT *const ref = shared_ref_data.data();
170+
DstT *const out = result.data();
171+
172+
cgh.single_task<Kernel<SrcT, NumElems, DstT, TestCaseT>>(
173+
[=]() SYCL_ESIMD_KERNEL {
174+
TestCaseT::template call_simd_ctor<SrcT, DstT, NumElems>(ref, out);
175+
});
176+
});
177+
queue.wait_and_throw();
178+
179+
for (size_t i = 0; i < result.size(); ++i) {
180+
// We ensure there is no UB here by preparing appropriate reference
181+
// values.
182+
const DstT &expected = static_cast<DstT>(ref_data[i]);
183+
const DstT &retrieved = result[i];
184+
if constexpr (type_traits::is_sycl_floating_point_v<DstT>) {
185+
// std::isnan() couldn't be called for integral types because it call is
186+
// ambiguous GitHub issue for that case:
187+
// https://github.com/microsoft/STL/issues/519
188+
if (!std::isnan(expected) || !std::isnan(retrieved)) {
189+
if (expected != retrieved) {
190+
// TODO add a function that will compare with defined accuracy
191+
// taking into account the possibility of UB on border values.
192+
// We don't have a such UB now because we are using 10f as maximum
193+
// value.
194+
if (!((expected + value<DstT>::pos_ulp(expected)) >= retrieved ||
195+
(expected - value<DstT>::pos_ulp(expected)) >= retrieved ||
196+
(expected + value<DstT>::pos_ulp(expected)) <= retrieved ||
197+
(expected - value<DstT>::pos_ulp(expected)) <= retrieved)) {
198+
passed = fail_test(i, retrieved, expected, src_data_type,
199+
dst_data_type);
200+
}
201+
}
202+
}
203+
} else {
204+
if (expected != retrieved) {
205+
passed =
206+
fail_test(i, retrieved, expected, src_data_type, dst_data_type);
207+
}
208+
}
209+
}
210+
211+
return passed;
212+
}
213+
214+
bool fail_test(size_t index, DstT retrieved, DstT expected,
215+
const std::string &src_data_type,
216+
const std::string &dst_data_type) {
217+
const auto description = ConvCtorTestDescription<DstT, NumElems, TestCaseT>(
218+
index, retrieved, expected, src_data_type, dst_data_type);
219+
log::fail(description);
220+
221+
return false;
222+
}
223+
};
224+
225+
} // namespace esimd_test::api::functional::ctors

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

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
//==------- ctor_converting_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+
// Test for simd converting constructor for core types.
18+
// This test uses different data types, dimensionality, base and step values and
19+
// different simd constructor invocation contexts.
20+
// The test do the following actions:
21+
// - construct simd with source data type, then construct simd with destination
22+
// type from the earlier constructed simd
23+
// - compare retrieved and expected values
24+
25+
#include "ctor_converting.hpp"
26+
27+
using namespace sycl::ext::intel::experimental::esimd;
28+
using namespace esimd_test::api::functional;
29+
30+
int main(int, char **) {
31+
sycl::queue queue(esimd_test::ESIMDSelector{},
32+
esimd_test::createExceptionHandler());
33+
34+
bool passed = true;
35+
36+
const auto fp_types = get_tested_types<tested_types::fp>();
37+
const auto uint_types = get_tested_types<tested_types::uint>();
38+
const auto sint_types = get_tested_types<tested_types::sint>();
39+
const auto core_types = get_tested_types<tested_types::core>();
40+
const auto single_size = get_sizes<1, 8>();
41+
const auto contexts =
42+
unnamed_type_pack<ctors::initializer, ctors::var_decl,
43+
ctors::rval_in_expr, ctors::const_ref>::generate();
44+
45+
// Run for specific combinations of types, vector length, base and step values
46+
// and invocation contexts.
47+
// The first types is the source types. the second types is the destination
48+
// types.
49+
passed &= for_all_combinations<ctors::run_test>(fp_types, single_size,
50+
fp_types, contexts, queue);
51+
passed &= for_all_combinations<ctors::run_test>(fp_types, single_size,
52+
uint_types, contexts, queue);
53+
passed &= for_all_combinations<ctors::run_test>(fp_types, single_size,
54+
sint_types, contexts, queue);
55+
passed &= for_all_combinations<ctors::run_test>(uint_types, single_size,
56+
core_types, contexts, queue);
57+
passed &= for_all_combinations<ctors::run_test>(sint_types, single_size,
58+
uint_types, contexts, queue);
59+
passed &= for_all_combinations<ctors::run_test>(sint_types, single_size,
60+
sint_types, contexts, queue);
61+
passed &= for_all_combinations<ctors::run_test>(sint_types, single_size,
62+
fp_types, contexts, queue);
63+
64+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
65+
return passed ? 0 : 1;
66+
}

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

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
//==------- ctor_converting_fp_extra.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 simd converting constructor for extra fp types.
18+
// This test uses extra fp data types with different dimensionality, base and
19+
// step values and different simd constructor invocation contexts.
20+
// The test do the following actions:
21+
// - construct simd with source data type, then construct simd with destination
22+
// type from the earlier constructed simd
23+
// - compare retrieved and expected values
24+
25+
#include "ctor_converting.hpp"
26+
27+
using namespace sycl::ext::intel::experimental::esimd;
28+
using namespace esimd_test::api::functional;
29+
30+
int main(int, char **) {
31+
sycl::queue queue(esimd_test::ESIMDSelector{},
32+
esimd_test::createExceptionHandler());
33+
34+
bool passed = true;
35+
36+
const auto fp_types = get_tested_types<tested_types::fp_extra>();
37+
const auto uint_types = get_tested_types<tested_types::uint>();
38+
const auto sint_types = get_tested_types<tested_types::sint>();
39+
const auto core_types = get_tested_types<tested_types::core>();
40+
const auto single_size = get_sizes<1, 8>();
41+
const auto contexts =
42+
unnamed_type_pack<ctors::initializer, ctors::var_decl,
43+
ctors::rval_in_expr, ctors::const_ref>::generate();
44+
45+
// Run for specific combinations of types, vector length, base and step values
46+
// and invocation contexts.
47+
// The first types is the source types. the second types is the destination
48+
// types.
49+
passed &= for_all_combinations<ctors::run_test>(fp_types, single_size,
50+
fp_types, contexts, queue);
51+
passed &= for_all_combinations<ctors::run_test>(fp_types, single_size,
52+
uint_types, contexts, queue);
53+
passed &= for_all_combinations<ctors::run_test>(fp_types, single_size,
54+
sint_types, contexts, queue);
55+
passed &= for_all_combinations<ctors::run_test>(uint_types, single_size,
56+
core_types, contexts, queue);
57+
passed &= for_all_combinations<ctors::run_test>(sint_types, single_size,
58+
uint_types, contexts, queue);
59+
passed &= for_all_combinations<ctors::run_test>(sint_types, single_size,
60+
sint_types, contexts, queue);
61+
passed &= for_all_combinations<ctors::run_test>(sint_types, single_size,
62+
fp_types, contexts, queue);
63+
64+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
65+
return passed ? 0 : 1;
66+
}

0 commit comments

Comments
 (0)