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

[SYCL] Extend the SpecConstants/vector-convolution-demo.cpp test #830

Merged
merged 4 commits into from Feb 22, 2022
Merged

[SYCL] Extend the SpecConstants/vector-convolution-demo.cpp test #830

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
816ae07
[SYCL] Extend the SpecConstants/vector-convolution-demo.cpp test
Feb 11, 2022
0b6c213
[SYCL] Fix compilation error in vector-convolution-demo.cpp
Feb 11, 2022
737b320
Add a second spec. constant with padding
Feb 15, 2022
3a4b86a
[NFC] Fix formatting issues
Feb 15, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
55 changes: 49 additions & 6 deletions SYCL/SpecConstants/2020/vector-convolution-demo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,18 +24,44 @@ struct coeff_struct_t {
std::array<std::array<float, 3>, 3> c;
};

coeff_t get_coefficients() {
return {{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}};
}
struct alignas(64) coeff_struct_aligned_t {
std::array<std::array<float, 3>, 3> c;
};

struct alignas(64) coeff_struct_aligned2_t {
std::array<std::array<float, 3>, 3> c;
int number;
};

coeff_struct_t get_coefficient_struct() {
template <typename T> constexpr T get_coefficients() {
return {{{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}}};
}

template <> constexpr coeff_t get_coefficients<coeff_t>() {
return {{{1.0, 2.0, 3.0}, {1.1, 2.1, 3.1}, {1.2, 2.2, 3.2}}};
}

constexpr specialization_id<coeff_t> coeff_id;

constexpr specialization_id<coeff_struct_t> coeff_struct_id;

// Represented in the IR as
// clang-format off
// { %struct.coeff_struct_aligned_t { %"class.std::array.0" zeroinitializer, [28 x i8] undef } }
// ~ padding ~
// clang-format on
constexpr specialization_id<coeff_struct_aligned_t> coeff_struct_aligned_id;
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It may make the test look a little odd, but would it make sense to have two of these? In that case we are sure that one will appear before another spec constant and we can be sure that the calculated offset isn't affected by the padding.

Copy link
Author

@ghost ghost Feb 11, 2022 •
edited by ghost
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@steffenlarsen Thank you for the comment. I'm sorry, maybe I've got your comment wrong. Should I add another spec constant with the same type as specialization_id<coeff_struct_aligned_t>, or just remove the constexpr specialization_id<coeff_struct_t> coeff_struct_id one because it has no padding and if the constant with padding has been processed well, the constant without padding should also be processed?

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My worry is that the changes to sycl-post-link could make it difficult for the runtime to correctly calculate the sizes and offsets of the specialization constants with padding. If the padded specialization constant is the last specialization constant we would likely not see this as there is no following specialization constant and as such the amount of padding does not matter to the runtime. Adding an extra specialization constant with specialization_id<coeff_struct_aligned_t> would ensure at least one of them isn't the last specialization constant, right?

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@steffenlarsen Thank you very much for the explanation, it's crystal clear now. I've added a new specialization constant after coeff_struct_aligned_id but I see locally the test failed (in fact, it failed even with one specialization constant containing padding). The test works!


// An extra specialization constant to check whether the runtime correctly
// calculates the sizes and offsets of the specialization constants with
// padding.
// Represented in the IR as
// clang-format off
// { %struct.coeff_struct_aligned2_t { %"class.std::array.0" zeroinitializer, i32 0, [24 x i8] undef } }
// ~ padding ~
// clang-format on
constexpr specialization_id<coeff_struct_aligned2_t> coeff_struct_aligned_id2;

template <typename IN>
float calc_conv(const coeff_t &coeff, const IN &in, item<2> item_id) {
float acc = 0;
Expand Down Expand Up @@ -64,8 +90,13 @@ void do_conv(buffer<float, 2> in, buffer<float, 2> out, CP coeff_provider) {

// Set the coefficient of the convolution as constant.
// This will build a specific kernel the coefficient available as literals.
cgh.set_specialization_constant<coeff_id>(get_coefficients());
cgh.set_specialization_constant<coeff_struct_id>(get_coefficient_struct());
cgh.set_specialization_constant<coeff_id>(get_coefficients<coeff_t>());
cgh.set_specialization_constant<coeff_struct_id>(
get_coefficients<coeff_struct_t>());
cgh.set_specialization_constant<coeff_struct_aligned_id>(
get_coefficients<coeff_struct_aligned_t>());
cgh.set_specialization_constant<coeff_struct_aligned_id2>(
get_coefficients<coeff_struct_aligned2_t>());
cgh.parallel_for<KernelName>(
in.get_range(), [=](item<2> item_id, kernel_handler h) {
auto coeff = coeff_provider(h);
Expand Down Expand Up @@ -126,6 +157,18 @@ int main() {

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

do_conv<class Convolution3>(input, output, [](kernel_handler &h) {
return h.get_specialization_constant<coeff_struct_aligned_id>().c;
});

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

do_conv<class Convolution4>(input, output, [](kernel_handler &h) {
return h.get_specialization_constant<coeff_struct_aligned_id2>().c;
});

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

std::cout << "Good computation!" << std::endl;
return 0;
}