Skip to content

Commit 8faef69

Browse files
yuriykochbb-sycl
yuriykoch
authored and
bb-sycl
committed
[SYCL][ESIMD] Rename tests on USM simd load constructors (intel#950)
* [SYCL][ESIMD] Rename tests on USM simd load constructors To distinguish from the simd load constructors from accessor Signed-off-by: Kochetkov, Yuriy <[email protected]>
1 parent 472c58e commit 8faef69

File tree

4 files changed

+202
-69
lines changed

4 files changed

+202
-69
lines changed

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

Lines changed: 9 additions & 64 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//===-- ctor_load.hpp - Functions for tests on simd load constructor definition.
1+
//===-- ctor_load.hpp - Functions for tests on simd load constructor.
22
// -------------------------------------------------------------------===//
33
//
44
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -8,14 +8,17 @@
88
//===----------------------------------------------------------------------===//
99
///
1010
/// \file
11-
/// This file provides functions for tests on simd load constructor.
11+
/// This file provides functions for tests on simd load constructor
1212
///
1313
//===----------------------------------------------------------------------===//
1414

1515
#pragma once
1616
<<<<<<< HEAD
1717
<<<<<<< HEAD
18+
<<<<<<< HEAD
1819
#define ESIMD_TESTS_DISABLE_DEPRECATED_TEST_DESCRIPTION_FOR_LOGS
20+
=======
21+
>>>>>>> 11cb2778d ([SYCL][ESIMD] Rename tests on USM simd load constructors (#950))
1922

2023
#include "common.hpp"
2124

@@ -35,72 +38,11 @@ namespace esimd = sycl::ext::intel::experimental::esimd;
3538

3639
namespace esimd_test::api::functional::ctors {
3740

38-
// Descriptor class for the case of calling constructor in initializer context.
39-
struct initializer {
40-
static std::string get_description() { return "initializer"; }
41-
42-
template <typename DataT, int NumElems, typename AlignmentT>
43-
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
44-
AlignmentT alignment) {
45-
esimd::simd<DataT, NumElems> simd_by_init =
46-
esimd::simd<DataT, NumElems>(ref_data, alignment);
47-
simd_by_init.copy_to(out);
48-
}
49-
};
50-
51-
// Descriptor class for the case of calling constructor in variable declaration
52-
// context.
53-
struct var_decl {
54-
static std::string get_description() { return "variable declaration"; }
55-
56-
template <typename DataT, int NumElems, typename AlignmentT>
57-
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
58-
AlignmentT alignment) {
59-
esimd::simd<DataT, NumElems> simd_by_var_decl(ref_data, alignment);
60-
simd_by_var_decl.copy_to(out);
61-
}
62-
};
63-
64-
// Descriptor class for the case of calling constructor in rvalue in an
65-
// expression context.
66-
struct rval_in_expr {
67-
static std::string get_description() { return "rvalue in an expression"; }
68-
69-
template <typename DataT, int NumElems, typename AlignmentT>
70-
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
71-
AlignmentT alignment) {
72-
esimd::simd<DataT, NumElems> simd_by_rval;
73-
simd_by_rval = esimd::simd<DataT, NumElems>(ref_data, alignment);
74-
simd_by_rval.copy_to(out);
75-
}
76-
};
77-
78-
// Descriptor class for the case of calling constructor in const reference
79-
// context.
80-
class const_ref {
81-
public:
82-
static std::string get_description() { return "const reference"; }
83-
84-
template <typename DataT, int NumElems, typename AlignmentT>
85-
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
86-
AlignmentT alignment) {
87-
call_simd_by_const_ref<DataT, NumElems>(
88-
esimd::simd<DataT, NumElems>(ref_data, alignment), out);
89-
}
90-
91-
private:
92-
template <typename DataT, int NumElems>
93-
static void
94-
call_simd_by_const_ref(const esimd::simd<DataT, NumElems> &simd_by_const_ref,
95-
DataT *out) {
96-
simd_by_const_ref.copy_to(out);
97-
}
98-
};
99-
10041
// Dummy kernel for submitting some code into device side.
10142
template <typename DataT, int NumElems, typename T, typename Alignment>
10243
struct Kernel_for_load_ctor;
10344

45+
// Alignment tags
10446
namespace alignment {
10547

10648
struct element {
@@ -188,6 +130,7 @@ class LoadCtorTestDescription : public ITestDescription {
188130
const std::string m_alignment_name;
189131
};
190132

133+
<<<<<<< HEAD
191134
<<<<<<< HEAD
192135
// The main test routine.
193136
// Using functor class to be able to iterate over the pre-defined data types.
@@ -354,4 +297,6 @@ class run_test {
354297
}
355298
};
356299

300+
=======
301+
>>>>>>> 11cb2778d ([SYCL][ESIMD] Rename tests on USM simd load constructors (#950))
357302
} // namespace esimd_test::api::functional::ctors

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

Lines changed: 184 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,184 @@
1+
//===-- ctor_load_usm.hpp - Functions for tests on USM simd load constructor.
2+
// -------------------------------------------------------------------===//
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 load constructor from USM
12+
/// pointer
13+
///
14+
//===----------------------------------------------------------------------===//
15+
16+
#pragma once
17+
#define ESIMD_TESTS_DISABLE_DEPRECATED_TEST_DESCRIPTION_FOR_LOGS
18+
19+
#include "ctor_load.hpp"
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, typename AlignmentT>
28+
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
29+
AlignmentT alignment) {
30+
esimd::simd<DataT, NumElems> simd_by_init =
31+
esimd::simd<DataT, NumElems>(ref_data, alignment);
32+
simd_by_init.copy_to(out);
33+
}
34+
};
35+
36+
// Descriptor class for the case of calling constructor in variable declaration
37+
// context.
38+
struct var_decl {
39+
static std::string get_description() { return "variable declaration"; }
40+
41+
template <typename DataT, int NumElems, typename AlignmentT>
42+
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
43+
AlignmentT alignment) {
44+
esimd::simd<DataT, NumElems> simd_by_var_decl(ref_data, alignment);
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_expr {
52+
static std::string get_description() { return "rvalue in an expression"; }
53+
54+
template <typename DataT, int NumElems, typename AlignmentT>
55+
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
56+
AlignmentT alignment) {
57+
esimd::simd<DataT, NumElems> simd_by_rval;
58+
simd_by_rval = esimd::simd<DataT, NumElems>(ref_data, alignment);
59+
simd_by_rval.copy_to(out);
60+
}
61+
};
62+
63+
// Descriptor class for the case of calling constructor in const reference
64+
// context.
65+
class const_ref {
66+
public:
67+
static std::string get_description() { return "const reference"; }
68+
69+
template <typename DataT, int NumElems, typename AlignmentT>
70+
static void call_simd_ctor(const DataT *ref_data, DataT *const out,
71+
AlignmentT alignment) {
72+
call_simd_by_const_ref<DataT, NumElems>(
73+
esimd::simd<DataT, NumElems>(ref_data, alignment), out);
74+
}
75+
76+
private:
77+
template <typename DataT, int NumElems>
78+
static void
79+
call_simd_by_const_ref(const esimd::simd<DataT, NumElems> &simd_by_const_ref,
80+
DataT *out) {
81+
simd_by_const_ref.copy_to(out);
82+
}
83+
};
84+
85+
// The main test routine.
86+
// Using functor class to be able to iterate over the pre-defined data types.
87+
template <typename DataT, typename SizeT, typename TestCaseT,
88+
typename AlignmentT>
89+
class run_test {
90+
static constexpr int NumElems = SizeT::value;
91+
using TestDescriptionT = LoadCtorTestDescription<NumElems, TestCaseT>;
92+
93+
public:
94+
bool operator()(sycl::queue &queue, const std::string &data_type,
95+
const std::string &alignment_name) {
96+
bool passed = true;
97+
log::trace<TestDescriptionT>(data_type, alignment_name);
98+
99+
const std::vector<DataT> ref_data = generate_ref_data<DataT, NumElems>();
100+
101+
// If current number of elements is equal to one, then run test with each
102+
// one value from reference data.
103+
// If current number of elements is greater than one, then run tests with
104+
// whole reference data.
105+
if constexpr (NumElems == 1) {
106+
for (size_t i = 0; i < ref_data.size(); ++i) {
107+
passed =
108+
run_verification(queue, {ref_data[i]}, data_type, alignment_name);
109+
}
110+
} else {
111+
passed = run_verification(queue, ref_data, data_type, alignment_name);
112+
}
113+
return passed;
114+
}
115+
116+
private:
117+
bool run_verification(sycl::queue &queue, const std::vector<DataT> &ref_data,
118+
const std::string &data_type,
119+
const std::string &alignment_name) {
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+
const size_t alignment_value =
126+
AlignmentT::template get_size<DataT, NumElems>();
127+
const size_t container_extra_size = alignment_value / sizeof(DataT) + 1;
128+
const size_t offset = 1;
129+
130+
shared_allocator<DataT> allocator(queue);
131+
shared_vector<DataT> result(NumElems, allocator);
132+
shared_vector<DataT> shared_ref_data(NumElems + container_extra_size +
133+
offset,
134+
shared_allocator<DataT>(queue));
135+
136+
const size_t object_size = NumElems * sizeof(DataT);
137+
size_t buffer_size = object_size + container_extra_size * sizeof(DataT);
138+
139+
// When we allocate USM there is a high probability that this memory will
140+
// have stronger alignment that required. We increment our pointer by fixed
141+
// offset value to avoid bigger alignment of USM shared.
142+
// The std::align can provide expected alignment on the small values of an
143+
// alignment.
144+
void *ref = shared_ref_data.data() + offset;
145+
if (std::align(alignment_value, object_size, ref, buffer_size) == nullptr) {
146+
return false;
147+
}
148+
DataT *const ref_aligned = static_cast<DataT *>(ref);
149+
150+
for (size_t i = 0; i < NumElems; ++i) {
151+
ref_aligned[i] = ref_data[i];
152+
}
153+
154+
queue.submit([&](sycl::handler &cgh) {
155+
DataT *const out = result.data();
156+
157+
cgh.single_task<
158+
Kernel_for_load_ctor<DataT, NumElems, TestCaseT, AlignmentT>>(
159+
[=]() SYCL_ESIMD_KERNEL {
160+
const auto alignment = AlignmentT::get_value();
161+
TestCaseT::template call_simd_ctor<DataT, NumElems>(ref_aligned,
162+
out, alignment);
163+
});
164+
});
165+
queue.wait_and_throw();
166+
167+
for (size_t i = 0; i < result.size(); ++i) {
168+
const auto &expected = ref_data[i];
169+
const auto &retrieved = result[i];
170+
171+
if (!are_bitwise_equal(expected, retrieved)) {
172+
passed = false;
173+
174+
log::fail(TestDescriptionT(data_type, alignment_name),
175+
"Unexpected value at index ", i, ", retrieved: ", retrieved,
176+
", expected: ", expected);
177+
}
178+
}
179+
180+
return passed;
181+
}
182+
};
183+
184+
} // namespace esimd_test::api::functional::ctors

‎SYCL/ESIMD/api/functional/ctors/ctor_load_core.cpp renamed to ‎SYCL/ESIMD/api/functional/ctors/ctor_load_usm_core.cpp

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//==------- ctor_load_core.cpp - DPC++ ESIMD on-device test ---------------==//
1+
//==------- ctor_load_usm_core.cpp - DPC++ ESIMD on-device test -----------==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -17,13 +17,17 @@
1717
// Test for simd load constructor.
1818
// The test uses reference data and different alignment flags. Invokes simd
1919
// constructors in different contexts with provided reference data and alignment
20-
// flag.
20+
// flag using USM pointer as input.
2121
// It is expected for destination simd instance to store a bitwise same data as
2222
// the reference one.
2323

24+
<<<<<<< HEAD:SYCL/ESIMD/api/functional/ctors/ctor_load_core.cpp
2425
<<<<<<< HEAD
2526
<<<<<<< HEAD
2627
#include "ctor_load.hpp"
28+
=======
29+
#include "ctor_load_usm.hpp"
30+
>>>>>>> 11cb2778d ([SYCL][ESIMD] Rename tests on USM simd load constructors (#950)):SYCL/ESIMD/api/functional/ctors/ctor_load_usm_core.cpp
2731

2832
using namespace esimd_test::api::functional;
2933

‎SYCL/ESIMD/api/functional/ctors/ctor_load_fp_extra.cpp renamed to ‎SYCL/ESIMD/api/functional/ctors/ctor_load_usm_fp_extra.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//==------- ctor_load_fp_extra.cpp - DPC++ ESIMD on-device test -----------==//
1+
//==------- ctor_load_usm_fp_extra.cpp - DPC++ ESIMD on-device test -------==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -17,11 +17,11 @@
1717
// Test for simd load constructor.
1818
// The test uses reference data and different alignment flags. Invokes simd
1919
// constructors in different contexts with provided reference data and alignment
20-
// flag.
20+
// flag using USM pointer as input.
2121
// It is expected for destination simd instance to store a bitwise same data as
2222
// the reference one.
2323

24-
#include "ctor_load.hpp"
24+
#include "ctor_load_usm.hpp"
2525

2626
using namespace esimd_test::api::functional;
2727

0 commit comments

Comments
 (0)