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

Commit 8c6a5ff

Browse files
vasilytricvasilytric
authored
[SYCL][ESIMD] Add test for esimd fill constructor (#598)
* [SYCL][ESIMD] Add test for simd fill constructor The test currently verifies the basic types; it does not yet check types requiring checks "if supported on device", such as "sycl::half". Also, the test does not yet verify 32-element vectors due to existing runtime problems in either device drivers or ESIMD implementation.
1 parent 0b6fc11 commit 8c6a5ff

File tree

5 files changed

+532
-15
lines changed

5 files changed

+532
-15
lines changed

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

Lines changed: 294 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,294 @@
1+
//===-- ctor_fill.hpp - Functions for tests on simd fill 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 fill constructor.
12+
///
13+
//===----------------------------------------------------------------------===//
14+
15+
#pragma once
16+
17+
#include "common.hpp"
18+
// For std::isnan
19+
#include <cmath>
20+
21+
namespace esimd = sycl::ext::intel::experimental::esimd;
22+
namespace esimd_functional = esimd_test::api::functional;
23+
24+
namespace esimd_test::api::functional::ctors {
25+
26+
// Descriptor class for the case of calling constructor in initializer context.
27+
struct initializer {
28+
static std::string get_description() { return "initializer"; }
29+
30+
template <typename DataT, int NumElems>
31+
static void call_simd_ctor(DataT init_value, DataT step, DataT *const out) {
32+
const auto simd_by_init = esimd::simd<DataT, NumElems>(init_value, step);
33+
simd_by_init.copy_to(out);
34+
}
35+
};
36+
37+
// Descriptor class for the case of calling constructor in variable declaration
38+
// context.
39+
struct var_dec {
40+
static std::string get_description() { return "variable declaration"; }
41+
42+
template <typename DataT, int NumElems>
43+
static void call_simd_ctor(DataT init_value, DataT step, DataT *const out) {
44+
const esimd::simd<DataT, NumElems> simd_by_var_decl(init_value, step);
45+
simd_by_var_decl.copy_to(out);
46+
}
47+
};
48+
49+
// Descriptor class for the case of calling constructor in rvalue in an
50+
// expression context.
51+
struct rval_in_express {
52+
static std::string get_description() { return "rvalue in an expression"; }
53+
54+
template <typename DataT, int NumElems>
55+
static void call_simd_ctor(DataT init_value, DataT step, DataT *const out) {
56+
esimd::simd<DataT, NumElems> simd_by_rval;
57+
simd_by_rval = esimd::simd<DataT, NumElems>(init_value, step);
58+
simd_by_rval.copy_to(out);
59+
}
60+
};
61+
62+
// Descriptor class for the case of calling constructor in const reference
63+
// context.
64+
class const_ref {
65+
public:
66+
static std::string get_description() { return "const reference"; }
67+
68+
template <typename DataT, int NumElems>
69+
static void call_simd_ctor(DataT init_value, DataT step, DataT *const out) {
70+
return call_simd_by_const_ref<DataT, NumElems>(
71+
esimd::simd<DataT, NumElems>(init_value, step), out);
72+
}
73+
74+
private:
75+
template <typename DataT, int NumElems>
76+
static void
77+
call_simd_by_const_ref(const esimd::simd<DataT, NumElems> &simd_by_const_ref,
78+
DataT *const out) {
79+
simd_by_const_ref.copy_to(out);
80+
}
81+
};
82+
83+
// Enumeration of possible values for base value and step that will be provided
84+
// into simd constructor.
85+
enum class init_val {
86+
min,
87+
max,
88+
zero,
89+
min_half,
90+
max_half,
91+
neg_inf,
92+
nan,
93+
positive,
94+
negative,
95+
denorm,
96+
inexact,
97+
ulp,
98+
ulp_half
99+
};
100+
101+
// Dummy kernel for submitting some code into device side.
102+
template <typename DataT, int NumElems, typename T, init_val BaseVal,
103+
init_val StepVal>
104+
struct kernel_for_fill;
105+
106+
// Constructing a value for step and base values that depends on input
107+
// parameters.
108+
template <typename DataT, init_val Value>
109+
DataT get_value(DataT base_val = DataT()) {
110+
if constexpr (Value == init_val::min) {
111+
return value<DataT>::lowest();
112+
} else if constexpr (Value == init_val::max) {
113+
return value<DataT>::max();
114+
} else if constexpr (Value == init_val::zero) {
115+
return 0;
116+
} else if constexpr (Value == init_val::positive) {
117+
return static_cast<DataT>(1.25);
118+
} else if constexpr (Value == init_val::negative) {
119+
return static_cast<DataT>(-10.75);
120+
} else if constexpr (Value == init_val::min_half) {
121+
return value<DataT>::lowest() / 2;
122+
} else if constexpr (Value == init_val::max_half) {
123+
return value<DataT>::max() / 2;
124+
} else if constexpr (Value == init_val::neg_inf) {
125+
return -value<DataT>::inf();
126+
} else if constexpr (Value == init_val::nan) {
127+
return value<DataT>::nan();
128+
} else if constexpr (Value == init_val::denorm) {
129+
return value<DataT>::denorm_min();
130+
} else if constexpr (Value == init_val::inexact) {
131+
return 0.1;
132+
} else if constexpr (Value == init_val::ulp) {
133+
return value<DataT>::pos_ulp(base_val);
134+
} else if constexpr (Value == init_val::ulp_half) {
135+
return value<DataT>::pos_ulp(base_val) / 2;
136+
} else {
137+
static_assert(Value != Value, "Unexpected enum value");
138+
}
139+
}
140+
141+
template <init_val Val> std::string init_val_to_string() {
142+
if constexpr (Val == init_val::min) {
143+
return "lowest";
144+
} else if constexpr (Val == init_val::max) {
145+
return "max";
146+
} else if constexpr (Val == init_val::zero) {
147+
return "zero";
148+
} else if constexpr (Val == init_val::positive) {
149+
return "positive";
150+
} else if constexpr (Val == init_val::negative) {
151+
return "negative";
152+
} else if constexpr (Val == init_val::min_half) {
153+
return "min_half";
154+
} else if constexpr (Val == init_val::max_half) {
155+
return "max_half";
156+
} else if constexpr (Val == init_val::neg_inf) {
157+
return "neg_inf";
158+
} else if constexpr (Val == init_val::nan) {
159+
return "nan";
160+
} else if constexpr (Val == init_val::denorm) {
161+
return "denorm";
162+
} else if constexpr (Val == init_val::inexact) {
163+
return "inexact";
164+
} else if constexpr (Val == init_val::ulp) {
165+
return "ulp";
166+
} else if constexpr (Val == init_val::ulp_half) {
167+
return "ulp_half";
168+
} else {
169+
static_assert(Val != Val, "Unexpected enum value");
170+
}
171+
}
172+
173+
template <typename DataT, int NumElems, typename ContextT, init_val BaseVal,
174+
init_val Step>
175+
class FillCtorTestDescription
176+
: public TestDescription<DataT, NumElems, ContextT> {
177+
public:
178+
FillCtorTestDescription(size_t index, DataT retrieved_val, DataT expected_val,
179+
const std::string &data_type)
180+
: TestDescription<DataT, NumElems, ContextT>(index, retrieved_val,
181+
expected_val, data_type) {}
182+
183+
std::string to_string() const override {
184+
std::string log_msg(
185+
TestDescription<DataT, NumElems, ContextT>::to_string());
186+
187+
log_msg += ", with base value: " + init_val_to_string<BaseVal>();
188+
log_msg += ", with step value: " + init_val_to_string<Step>();
189+
190+
return log_msg;
191+
}
192+
};
193+
194+
template <typename DataT, int NumElems, typename TestCaseT, typename BaseVal,
195+
typename Step>
196+
class run_test {
197+
public:
198+
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+
204+
shared_vector<DataT> result(NumElems, shared_allocator<DataT>(queue));
205+
206+
const auto base_value = get_value<DataT, BaseVal::value>();
207+
const auto step_value = get_value<DataT, Step::value>(base_value);
208+
209+
queue.submit([&](sycl::handler &cgh) {
210+
DataT *const out = result.data();
211+
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+
});
218+
});
219+
queue.wait_and_throw();
220+
bool passed = true;
221+
222+
// Verify the base value was passed as-is
223+
if (!are_bitwise_equal(result[0], base_value)) {
224+
passed = fail_test(0, result[0], base_value, data_type);
225+
}
226+
227+
// Verify the step value works as expected being passed to the fill
228+
// constructor.
229+
DataT expected_value = base_value;
230+
for (size_t i = 1; i < result.size(); ++i) {
231+
if constexpr (BaseVal::value == init_val::nan ||
232+
Step::value == init_val::nan) {
233+
234+
if (!std::isnan(result[i])) {
235+
passed = false;
236+
237+
// TODO: Make ITestDescription architecture more flexible.
238+
// We are assuming that the NaN opcode may differ
239+
std::string log_msg = "Failed for simd<";
240+
log_msg += data_type + ", " + std::to_string(NumElems) + ">";
241+
log_msg += ", with context: " + TestCaseT::get_description();
242+
log_msg += ". The element at index: " + std::to_string(i) +
243+
", 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>();
247+
248+
log::note(log_msg);
249+
}
250+
} else {
251+
252+
expected_value += step_value;
253+
if (!are_bitwise_equal(result[i], expected_value)) {
254+
passed = fail_test(i, result[i], expected_value, data_type);
255+
}
256+
}
257+
}
258+
return passed;
259+
}
260+
261+
private:
262+
bool fail_test(size_t index, DataT retrieved, DataT expected,
263+
const std::string &data_type) {
264+
const auto description =
265+
FillCtorTestDescription<DataT, NumElems, TestCaseT, BaseVal::value,
266+
Step::value>(index, retrieved, expected,
267+
data_type);
268+
log::fail(description);
269+
270+
return false;
271+
}
272+
};
273+
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+
294+
} // namespace esimd_test::api::functional::ctors

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

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
//==------- ctor_fill_accuracy_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 simd<float, 32> fills with unexpected values while base value is denorm
18+
// and step is ulp. The SIMD_RUN_TEST_WITH_VECTOR_LEN_32 macros must be enabled
19+
// when it is resolved.
20+
//
21+
// The test verifies that simd fill constructor has no precision differences.
22+
// The test do the following actions:
23+
// - call simd with predefined base and step values
24+
// - bitwise comparing that output[0] value is equal to base value and
25+
// output[i] is equal to output[i -1] + step_value
26+
27+
#include "ctor_fill.hpp"
28+
29+
using namespace sycl::ext::intel::experimental::esimd;
30+
using namespace esimd_test::api::functional;
31+
32+
int main(int, char **) {
33+
sycl::queue queue(esimd_test::ESIMDSelector{},
34+
esimd_test::createExceptionHandler());
35+
36+
bool passed = true;
37+
38+
const auto fp_types = get_tested_types<tested_types::fp>();
39+
const auto single_dim = values_pack<8>();
40+
41+
// Run for specific combinations of types, base and step values and vector
42+
// length.
43+
// The first init_val value it's a base value and the second init_val value
44+
// it's a step value.
45+
#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);
49+
#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);
63+
64+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
65+
return passed ? 0 : 1;
66+
}

0 commit comments

Comments
 (0)