Skip to content

Commit bfbb9c0

Browse files
vasilytricbb-sycl
vasilytric
authored and
bb-sycl
committed
[SYCL][ESIMD] Add tests for simd broadcast constructor (intel#690)
* [SYCL][ESIMD] Add test on simd broadcast constructor
1 parent 0757766 commit bfbb9c0

File tree

7 files changed

+380
-14
lines changed

7 files changed

+380
-14
lines changed

‎SYCL/ESIMD/api/functional/common.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -47,7 +47,7 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
4747
} // namespace details
4848

4949
// Class used as a kernel ID.
50-
template <typename DataT, int NumElems, typename T> struct Kernel;
50+
template <typename DataT, int NumElems, typename...> struct Kernel;
5151

5252
template <typename DataT>
5353
using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;

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

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

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

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,80 @@
1+
//==------- ctor_copy.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+
// XRUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
15+
// XRUN: %GPU_RUN_PLACEHOLDER %t.out
16+
// RUN: false
17+
// XFAIL: *
18+
// TODO The simd filled with unexpected values.
19+
//
20+
// Test for simd broadcast constructor.
21+
// This test uses different data types, sizes and different simd constructor
22+
// invocation contexts.
23+
// Type of a value that will be provided to the broadcast constructor may be
24+
// differ, than value, that will be provided to the simd when it will be
25+
// constructed. It is expected for a new simd instance to store same data as the
26+
// one passed as the source simd constructor.
27+
28+
#include "ctor_broadcast.hpp"
29+
30+
using namespace sycl::ext::intel::experimental::esimd;
31+
using namespace esimd_test::api::functional;
32+
33+
int main(int, char **) {
34+
sycl::queue queue(esimd_test::ESIMDSelector{},
35+
esimd_test::createExceptionHandler());
36+
37+
bool passed = true;
38+
39+
const auto uint_types = get_tested_types<tested_types::uint>();
40+
const auto sint_types = get_tested_types<tested_types::sint>();
41+
const auto fp_types = get_tested_types<tested_types::fp>();
42+
const auto core_types = get_tested_types<tested_types::core>();
43+
const auto int_type = named_type_pack<int>::generate("int");
44+
using use_positive_value_only = std::true_type;
45+
using use_ref_conv_values = std::false_type;
46+
const auto single_size = get_sizes<8>();
47+
const auto two_sizes = get_sizes<1, 8>();
48+
const auto all_sizes = get_all_sizes();
49+
const auto all_contexts =
50+
unnamed_type_pack<ctors::initializer, ctors::var_decl,
51+
ctors::rval_in_expr, ctors::const_ref>::generate();
52+
const auto context = unnamed_type_pack<ctors::var_decl>::generate();
53+
54+
// Run for specific combinations of types, vector length, base and step values
55+
// and invocation contexts.
56+
// The source types is the first types, that provided to the
57+
// "for_all_combinations" the destination types is the second types that
58+
// provided to the "for_all_combinations".
59+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
60+
int_type, two_sizes, uint_types, all_contexts, queue);
61+
passed &= for_all_combinations<ctors::run_test_with_positive_value_only>(
62+
core_types, all_sizes, core_types, all_contexts, queue);
63+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
64+
fp_types, single_size, fp_types, context, queue);
65+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
66+
fp_types, single_size, uint_types, context, queue);
67+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
68+
fp_types, single_size, sint_types, context, queue);
69+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
70+
uint_types, single_size, core_types, context, queue);
71+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
72+
sint_types, single_size, uint_types, context, queue);
73+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
74+
sint_types, single_size, sint_types, context, queue);
75+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
76+
sint_types, single_size, fp_types, context, queue);
77+
78+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
79+
return passed ? 0 : 1;
80+
}

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

Lines changed: 59 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,59 @@
1+
//==------- ctor_broadcast_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+
// XFAIL: *
17+
// TODO The simd filled with unexpected values if the source type is floating
18+
// point type.
19+
//
20+
// Test for simd broadcast constructor.
21+
// This test uses fp extra data types, sizes and different simd constructor
22+
// invocation contexts.
23+
// Type of a value that will be provided to the broadcast constructor may be
24+
// differ, than value, that will be provided to the simd when it will be
25+
// constructed. It is expected for a new simd instance to store same data as the
26+
// one passed as the source simd constructor.
27+
28+
#include "ctor_broadcast.hpp"
29+
30+
using namespace sycl::ext::intel::experimental::esimd;
31+
using namespace esimd_test::api::functional;
32+
33+
int main(int, char **) {
34+
sycl::queue queue(esimd_test::ESIMDSelector{},
35+
esimd_test::createExceptionHandler());
36+
37+
bool passed = true;
38+
39+
const auto uint_types = get_tested_types<tested_types::uint>();
40+
const auto sint_types = get_tested_types<tested_types::sint>();
41+
const auto fp_extra_types = get_tested_types<tested_types::fp_extra>();
42+
const auto single_size = get_sizes<8>();
43+
const auto context = unnamed_type_pack<ctors::var_decl>::generate();
44+
45+
// Run for specific combinations of types, vector length, base and step values
46+
// and invocation contexts.
47+
// The source types is the first types, that provided to the
48+
// "for_all_combinations" the destination types is the second types that
49+
// provided to the "for_all_combinations".
50+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
51+
fp_extra_types, single_size, fp_extra_types, context, queue);
52+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
53+
fp_extra_types, single_size, uint_types, context, queue);
54+
passed &= for_all_combinations<ctors::run_test_with_all_values>(
55+
fp_extra_types, single_size, sint_types, context, queue);
56+
57+
std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
58+
return passed ? 0 : 1;
59+
}

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,8 +24,8 @@ namespace esimd_test::api::functional::operators {
2424

2525
// The main test routine.
2626
// Using functor class to be able to iterate over the pre-defined data types.
27-
template <typename DataT, typename DimT, typename TestCaseT> class run_test {
28-
static constexpr int NumElems = DimT::value;
27+
template <typename DataT, typename SizeT, typename TestCaseT> class run_test {
28+
static constexpr int NumElems = SizeT::value;
2929

3030
public:
3131
bool operator()(sycl::queue &queue, const std::string &data_type) {

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -63,14 +63,14 @@ int main(int, char **) {
6363

6464
bool passed = true;
6565

66-
const auto types = get_tested_types<tested_types::core>();
67-
const auto dims = get_all_dimensions();
66+
const auto types = get_tested_types<tested_types::all>();
67+
const auto all_sizes = get_all_sizes();
6868

6969
const auto context =
7070
unnamed_type_pack<move_assignment, copy_assignment>::generate();
7171

72-
passed &=
73-
for_all_combinations<operators::run_test>(types, dims, context, queue);
72+
passed &= for_all_combinations<operators::run_test>(types, all_sizes, context,
73+
queue);
7474

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

0 commit comments

Comments
 (0)