Skip to content

Commit

Permalink
Test more algorithms
Browse files Browse the repository at this point in the history
  • Loading branch information
victor-eds committed Feb 6, 2024
1 parent 9dab85a commit a143612
Show file tree
Hide file tree
Showing 3 changed files with 28 additions and 10 deletions.
Original file line number Diff line number Diff line change
@@ -0,0 +1,8 @@
// RUN: %{build} -fsycl-embed-ir -o %t.out
// RUN: %{run} %t.out

#include "./reduction.hpp"

int main() {
test<detail::reduction::strategy::group_reduce_and_atomic_cross_wg>();
}
Original file line number Diff line number Diff line change
@@ -0,0 +1,8 @@
// RUN: %{build} -fsycl-embed-ir -o %t.out
// RUN: %{run} %t.out

#include "./reduction.hpp"

int main() {
test<detail::reduction::strategy::local_mem_tree_and_atomic_cross_wg>();
}
22 changes: 12 additions & 10 deletions sycl/test-e2e/KernelFusion/Reduction/reduction.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -8,13 +8,11 @@

using namespace sycl;

constexpr inline size_t globalSize = 512;

template <detail::reduction::strategy Strategy> struct is_fusion_supported {
constexpr static inline bool value =
detail::reduction::strategy::group_reduce_and_last_wg_detection <=
Strategy &&
Strategy < detail::reduction::strategy::group_reduce_and_atomic_cross_wg;
Strategy < detail::reduction::strategy::group_reduce_and_multiple_kernels;
};

template <detail::reduction::strategy Strategy>
Expand All @@ -25,7 +23,9 @@ template <detail::reduction::strategy Strategy, bool Fuse>
void test(nd_range<1> ndr) {
static_assert(is_fusion_supported_v<Strategy>,
"Testing unsupported algorithm");
std::array<int, globalSize> data;

auto globalSize = static_cast<int>(ndr.get_global_range().size());
std::vector<int> data(globalSize);
int sumRes = 0;
int maxRes = 0;

Expand Down Expand Up @@ -70,17 +70,19 @@ void test(nd_range<1> ndr) {
}
}

constexpr int expectedMax = globalSize - 1;
constexpr int expectedSum = globalSize * expectedMax / 2;
int expectedMax = globalSize - 1;
int expectedSum = globalSize * expectedMax / 2;

assert(sumRes == expectedSum);
assert(maxRes == expectedMax);
}

template <detail::reduction::strategy Strategy> void test() {
for (size_t localSize = 1; localSize <= globalSize; localSize *= 2) {
nd_range<1> ndr{globalSize, localSize};
test<Strategy, true>(ndr);
test<Strategy, false>(ndr);
for (size_t globalSize : {16, 512}) {
for (size_t localSize = 1; localSize <= globalSize; localSize *= 2) {
nd_range<1> ndr{globalSize, localSize};
test<Strategy, true>(ndr);
test<Strategy, false>(ndr);
}
}
}

0 comments on commit a143612

Please sign in to comment.