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

Commit c1366f1

Browse files
vasilytricvasilytric
authored
[SYCL][ESIMD] Split tests on simd constructors into core and fp_extra (#748)
* [SYCL][ESIMD] Split tests into core and fp_extra
1 parent 1548e68 commit c1366f1

20 files changed

+898
-404
lines changed

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

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

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

Lines changed: 6 additions & 145 deletions
Original file line numberDiff line numberDiff line change
@@ -22,162 +22,23 @@
2222
// - use std::move() to provide it to simd constructor
2323
// - bitwise compare expected and retrieved values
2424

25-
#include "common.hpp"
25+
#include "ctor_array.hpp"
2626

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

16929
int main(int, char **) {
17030
sycl::queue queue(esimd_test::ESIMDSelector{},
17131
esimd_test::createExceptionHandler());
17232

17333
bool passed = true;
17434

175-
const auto types = get_tested_types<tested_types::all>();
35+
const auto types = get_tested_types<tested_types::core>();
17636
const auto dims = get_all_dimensions();
177-
const auto contexts = unnamed_type_pack<initializer, var_decl, rval_in_expr,
178-
const_ref>::generate();
37+
const auto contexts =
38+
unnamed_type_pack<ctors::initializer, ctors::var_decl,
39+
ctors::rval_in_expr, ctors::const_ref>::generate();
17940

180-
passed &= for_all_combinations<run_test>(types, dims, contexts, queue);
41+
passed &= for_all_combinations<ctors::run_test>(types, dims, contexts, queue);
18142

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

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
//==------- ctor_array_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 constructor from an array.
18+
// This test uses extra fp data types, dimensionality and different simd
19+
// constructor invocation contexts.
20+
// The test does the following actions:
21+
// - construct fixed-size array that filled with reference values
22+
// - use std::move() to provide it to simd constructor
23+
// - bitwise compare expected and retrieved values
24+
25+
#include "ctor_array.hpp"
26+
27+
using namespace esimd_test::api::functional;
28+
29+
int main(int, char **) {
30+
sycl::queue queue(esimd_test::ESIMDSelector{},
31+
esimd_test::createExceptionHandler());
32+
33+
bool passed = true;
34+
35+
const auto types = get_tested_types<tested_types::fp_extra>();
36+
const auto dims = get_all_dimensions();
37+
const auto contexts =
38+
unnamed_type_pack<ctors::initializer, ctors::var_decl,
39+
ctors::rval_in_expr, ctors::const_ref>::generate();
40+
41+
passed &= for_all_combinations<ctors::run_test>(types, dims, contexts, queue);
42+
43+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
44+
return passed ? 0 : 1;
45+
}

0 commit comments

Comments
 (0)