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

Commit 816ae07

Browse files
author
Pavel Samolysov
committed
[SYCL] Extend the SpecConstants/vector-convolution-demo.cpp test
The test is extended with a specialization constant built upon an aligned structure. The CFE generates a padding represented as '[size x i8] undef' for such structures and the sycl-post-link tool crashed during property set generation for the specialization constant. The bug has been fixed in intel/llvm#5538, this extension is the E2E test for the fix to ensure that the default values and offsets given to the runtime are correct. Signed-off-by: Pavel Samolysov <[email protected]>
1 parent cadc45a commit 816ae07

File tree

1 file changed

+23
-0
lines changed

1 file changed

+23
-0
lines changed

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

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

27+
struct alignas(64) coeff_struct_aligned_t {
28+
std::array<std::array<float, 3>, 3> c;
29+
};
30+
2731
coeff_t get_coefficients() {
2832
return {{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}};
2933
}
@@ -32,10 +36,21 @@ coeff_struct_t get_coefficient_struct() {
3236
return {{{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}}};
3337
}
3438

39+
// represented in the IR as
40+
// clang-format off
41+
// { %struct.coeff_struct_aligned_t { %"class.std::array.0" zeroinitializer, [28 x i8] undef } }
42+
// ~ padding ~
43+
// clang-format on
44+
coeff_struct_aligned_t get_coefficient_struct_aligned() {
45+
return {{{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}}};
46+
}
47+
3548
constexpr specialization_id<coeff_t> coeff_id;
3649

3750
constexpr specialization_id<coeff_struct_t> coeff_struct_id;
3851

52+
constexpr specialization_id<coeff_struct_aligned_t> coeff_struct_aligned_id;
53+
3954
template <typename IN>
4055
float calc_conv(const coeff_t &coeff, const IN &in, item<2> item_id) {
4156
float acc = 0;
@@ -66,6 +81,8 @@ void do_conv(buffer<float, 2> in, buffer<float, 2> out, CP coeff_provider) {
6681
// This will build a specific kernel the coefficient available as literals.
6782
cgh.set_specialization_constant<coeff_id>(get_coefficients());
6883
cgh.set_specialization_constant<coeff_struct_id>(get_coefficient_struct());
84+
cgh.set_specialization_constant<coeff_struct_aligned_id>(
85+
get_coefficient_aligned_struct());
6986
cgh.parallel_for<KernelName>(
7087
in.get_range(), [=](item<2> item_id, kernel_handler h) {
7188
auto coeff = coeff_provider(h);
@@ -126,6 +143,12 @@ int main() {
126143

127144
compare_result(host_accessor{output, read_only}, expected);
128145

146+
do_conv<class Convolution3>(input, output, [](kernel_handler &h) {
147+
return h.get_specialization_constant<coeff_struct_aligned_id>().c;
148+
});
149+
150+
compare_result(host_accessor{output, read_only}, expected);
151+
129152
std::cout << "Good computation!" << std::endl;
130153
return 0;
131154
}

0 commit comments

Comments
 (0)