Skip to content

Commit 520c35e

Browse files
Pavel Samolysovbb-sycl
authored andcommitted
[SYCL] Add a vector convolution demo of using specialization constants (intel#747)
The code is based on the example of using specialization constants from the text of the SYCL 2020 specification, see https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_example_usage Signed-off-by: Pavel Samolysov <[email protected]>
1 parent bc7f886 commit 520c35e

File tree

1 file changed

+20
-0
lines changed

1 file changed

+20
-0
lines changed

SYCL/SpecConstants/2020/vector-convolution-demo.cpp

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@ struct coeff_struct_t {
2424
std::array<std::array<float, 3>, 3> c;
2525
};
2626

27+
<<<<<<< HEAD
2728
struct alignas(64) coeff_struct_aligned_t {
2829
std::array<std::array<float, 3>, 3> c;
2930
};
@@ -39,12 +40,21 @@ template <typename T> constexpr T get_coefficients() {
3940

4041
template <> constexpr coeff_t get_coefficients<coeff_t>() {
4142
return {{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}};
43+
=======
44+
coeff_t get_coefficients() {
45+
return {{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}};
46+
}
47+
48+
coeff_struct_t get_coefficient_struct() {
49+
return {{{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}}};
50+
>>>>>>> 6c7a8e6f7 ([SYCL] Add a vector convolution demo of using specialization constants (#747))
4251
}
4352

4453
constexpr specialization_id<coeff_t> coeff_id;
4554

4655
constexpr specialization_id<coeff_struct_t> coeff_struct_id;
4756

57+
<<<<<<< HEAD
4858
// Represented in the IR as
4959
// clang-format off
5060
// { %struct.coeff_struct_aligned_t { %"class.std::array.0" zeroinitializer, [28 x i8] undef } }
@@ -62,6 +72,8 @@ constexpr specialization_id<coeff_struct_aligned_t> coeff_struct_aligned_id;
6272
// clang-format on
6373
constexpr specialization_id<coeff_struct_aligned2_t> coeff_struct_aligned_id2;
6474

75+
=======
76+
>>>>>>> 6c7a8e6f7 ([SYCL] Add a vector convolution demo of using specialization constants (#747))
6577
template <typename IN>
6678
float calc_conv(const coeff_t &coeff, const IN &in, item<2> item_id) {
6779
float acc = 0;
@@ -90,13 +102,18 @@ void do_conv(buffer<float, 2> in, buffer<float, 2> out, CP coeff_provider) {
90102

91103
// Set the coefficient of the convolution as constant.
92104
// This will build a specific kernel the coefficient available as literals.
105+
<<<<<<< HEAD
93106
cgh.set_specialization_constant<coeff_id>(get_coefficients<coeff_t>());
94107
cgh.set_specialization_constant<coeff_struct_id>(
95108
get_coefficients<coeff_struct_t>());
96109
cgh.set_specialization_constant<coeff_struct_aligned_id>(
97110
get_coefficients<coeff_struct_aligned_t>());
98111
cgh.set_specialization_constant<coeff_struct_aligned_id2>(
99112
get_coefficients<coeff_struct_aligned2_t>());
113+
=======
114+
cgh.set_specialization_constant<coeff_id>(get_coefficients());
115+
cgh.set_specialization_constant<coeff_struct_id>(get_coefficient_struct());
116+
>>>>>>> 6c7a8e6f7 ([SYCL] Add a vector convolution demo of using specialization constants (#747))
100117
cgh.parallel_for<KernelName>(
101118
in.get_range(), [=](item<2> item_id, kernel_handler h) {
102119
auto coeff = coeff_provider(h);
@@ -157,6 +174,7 @@ int main() {
157174

158175
compare_result(host_accessor{output, read_only}, expected);
159176

177+
<<<<<<< HEAD
160178
do_conv<class Convolution3>(input, output, [](kernel_handler &h) {
161179
return h.get_specialization_constant<coeff_struct_aligned_id>().c;
162180
});
@@ -169,6 +187,8 @@ int main() {
169187

170188
compare_result(host_accessor{output, read_only}, expected);
171189

190+
=======
191+
>>>>>>> 6c7a8e6f7 ([SYCL] Add a vector convolution demo of using specialization constants (#747))
172192
std::cout << "Good computation!" << std::endl;
173193
return 0;
174194
}

0 commit comments

Comments
 (0)