From 591bd92f7d7c2494192690a6b56bc1b601f86119 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Mon, 24 May 2021 09:53:57 +0300 Subject: [PATCH 01/16] [SYCL] Added test for constexpr half operations Signed-off-by: mdimakov --- SYCL/Basic/half_type.cpp | 37 +++++++++++++++++++++++++++++++++++++ 1 file changed, 37 insertions(+) diff --git a/SYCL/Basic/half_type.cpp b/SYCL/Basic/half_type.cpp index 464f9ba627..7465d45e3c 100644 --- a/SYCL/Basic/half_type.cpp +++ b/SYCL/Basic/half_type.cpp @@ -192,7 +192,44 @@ inline bool bitwise_comparison_fp32(const half val, const uint32_t exp) { return reinterpret_cast(fp32) == exp; } +constexpr void constexpr_verify_add() { + constexpr half a{5.0}, b{2.0}, ref{7.0}; + constexpr half result = a + b; + constexpr half diff = result - ref; + static_assert(std::fabs(static_cast(diff)) < + std::numeric_limits::epsilon(), "Constexpr add is wrong"); +} + +constexpr void constexpr_verify_sub() { + constexpr half a{5.0f}, b{2.0}, ref{3.0}; + constexpr half result = a - b; + constexpr half diff = result - ref; + static_assert(std::fabs(static_cast(diff)) < + std::numeric_limits::epsilon(), "Constexpr sub is wrong"); +} + +constexpr void constexpr_verify_mul() { + constexpr half a{5.0f}, b{2.0}, ref{10.0}; + constexpr half result = a * b; + constexpr half diff = result - ref; + static_assert(std::fabs(static_cast(diff)) < + std::numeric_limits::epsilon(), "Constexpr mul is wrong"); +} + +constexpr void constexpr_verify_div() { + constexpr half a{5.0f}, b{2.0}, ref{2.5}; + constexpr half result = a / b; + constexpr half diff = result - ref; + static_assert(std::fabs(static_cast(diff)) < + std::numeric_limits::epsilon(), "Constexpr div is wrong"); +} + int main() { + constexpr_verify_add(); + constexpr_verify_sub(); + constexpr_verify_mul(); + constexpr_verify_div(); + device dev{default_selector()}; if (!dev.is_host() && !dev.has_extension("cl_khr_fp16")) { std::cout << "This device doesn't support the extension cl_khr_fp16" From 1d07345bc57955676e66ac9d98f9c9421cb1da71 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Fri, 4 Jun 2021 13:05:25 +0300 Subject: [PATCH 02/16] [SYCL] Test for sycl function objects --- .../exclusive_scan_sycl2020.cpp | 162 +++++++++++++++++ .../inclusive_scan_sycl2020.cpp | 163 ++++++++++++++++++ SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 98 +++++++++++ 3 files changed, 423 insertions(+) create mode 100644 SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp create mode 100644 SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp create mode 100644 SYCL/GroupAlgorithm/reduce_sycl2020.cpp diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp new file mode 100644 index 0000000000..1987e52656 --- /dev/null +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -0,0 +1,162 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3. +// That requires either adding a switch to clang (-spirv-max-version=1.3) or +// raising the spirv version from 1.1. to 1.3 for spirv translator +// unconditionally. Using operators specific for spirv 1.3 and higher with +// -spirv-max-version=1.1 being set by default causes assert/check fails +// in spirv translator. +// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ + %t13.out + +#include "support.h" +#include +#include +#include +#include +#include +#include +using namespace sycl; + +template +class exclusive_scan_kernel; + +// std::exclusive_scan isn't implemented yet, so use serial implementation +// instead +namespace emu { +template +OutputIterator exclusive_scan(InputIterator first, InputIterator last, + OutputIterator result, T init, + BinaryOperation binary_op) { + T partial = init; + for (InputIterator it = first; it != last; ++it) { + *(result++) = partial; + partial = binary_op(partial, *it); + } + return result; +} +} // namespace emu + +template +void test(queue q, InputContainer input, OutputContainer output, + BinaryOperation binary_op, + typename OutputContainer::value_type identity) { + typedef typename InputContainer::value_type InputT; + typedef typename OutputContainer::value_type OutputT; + typedef class exclusive_scan_kernel kernel_name0; + typedef class exclusive_scan_kernel kernel_name1; + typedef class exclusive_scan_kernel kernel_name2; + typedef class exclusive_scan_kernel kernel_name3; + OutputT init = 42; + size_t N = input.size(); + size_t G = 64; + std::vector expected(N); + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + int lid = it.get_local_id(0); + out[lid] = exclusive_scan_over_group(g, in[lid], binary_op); + }); + }); + } + emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), + identity, binary_op); + assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + int lid = it.get_local_id(0); + out[lid] = exclusive_scan_over_group(g, in[lid], init, binary_op); + }); + }); + } + emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init, + binary_op); + assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + joint_exclusive_scan(g, in.get_pointer(), in.get_pointer() + N, + out.get_pointer(), binary_op); + }); + }); + } + emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), + identity, binary_op); + assert(std::equal(output.begin(), output.begin() + N, expected.begin())); + + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + joint_exclusive_scan(g, in.get_pointer(), in.get_pointer() + N, + out.get_pointer(), init, binary_op); + }); + }); + } + emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init, + binary_op); + assert(std::equal(output.begin(), output.begin() + N, expected.begin())); +} + +int main() { + queue q; + if (!isSupportedDevice(q.get_device())) { + std::cout << "Skipping test\n"; + return 0; + } + + constexpr int N = 128; + std::array input; + std::array output; + std::iota(input.begin(), input.end(), 0); + std::fill(output.begin(), output.end(), 0); + + test(q, input, output, sycl::plus<>(), 0); + test(q, input, output, sycl::minimum<>(), + std::numeric_limits::max()); + test(q, input, output, sycl::maximum<>(), + std::numeric_limits::lowest()); + + test(q, input, output, sycl::plus(), 0); + test(q, input, output, sycl::minimum(), + std::numeric_limits::max()); + test(q, input, output, sycl::maximum(), + std::numeric_limits::lowest()); + +#ifdef SPIRV_1_3 + test(q, input, output, sycl::multiplies(), + 1); + test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), ~0); +#endif // SPIRV_1_3 + + std::cout << "Test passed." << std::endl; +} diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp new file mode 100644 index 0000000000..f855c0717e --- /dev/null +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -0,0 +1,163 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3. +// That requires either adding a switch to clang (-spirv-max-version=1.3) or +// raising the spirv version from 1.1. to 1.3 for spirv translator +// unconditionally. Using operators specific for spirv 1.3 and higher with +// -spirv-max-version=1.1 being set by default causes assert/check fails +// in spirv translator. +// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ + %t13.out + +#include "support.h" +#include +#include +#include +#include +#include +#include +using namespace sycl; + +template +class inclusive_scan_kernel; + +// std::inclusive_scan isn't implemented yet, so use serial implementation +// instead +namespace emu { +template +OutputIterator inclusive_scan(InputIterator first, InputIterator last, + OutputIterator result, BinaryOperation binary_op, + T init) { + T partial = init; + for (InputIterator it = first; it != last; ++it) { + partial = binary_op(partial, *it); + *(result++) = partial; + } + return result; +} +} // namespace emu + +template +void test(queue q, InputContainer input, OutputContainer output, + BinaryOperation binary_op, + typename OutputContainer::value_type identity) { + typedef typename InputContainer::value_type InputT; + typedef typename OutputContainer::value_type OutputT; + typedef class inclusive_scan_kernel kernel_name0; + typedef class inclusive_scan_kernel kernel_name1; + typedef class inclusive_scan_kernel kernel_name2; + typedef class inclusive_scan_kernel kernel_name3; + OutputT init = 42; + size_t N = input.size(); + size_t G = 64; + std::vector expected(N); + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + int lid = it.get_local_id(0); + out[lid] = inclusive_scan_over_group(g, in[lid], binary_op); + }); + }); + } + emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), + binary_op, identity); + assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + int lid = it.get_local_id(0); + out[lid] = inclusive_scan_over_group(g, in[lid], binary_op, init); + }); + }); + } + emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), + binary_op, init); + assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + joint_inclusive_scan(g, in.get_pointer(), in.get_pointer() + N, + out.get_pointer(), binary_op); + }); + }); + } + emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), + binary_op, identity); + assert(std::equal(output.begin(), output.begin() + N, expected.begin())); + + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for(nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + joint_inclusive_scan(g, in.get_pointer(), in.get_pointer() + N, + out.get_pointer(), binary_op, init); + }); + }); + } + emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), + binary_op, init); + assert(std::equal(output.begin(), output.begin() + N, expected.begin())); +} + +int main() { + queue q; + if (!isSupportedDevice(q.get_device())) { + std::cout << "Skipping test\n"; + return 0; + } + + constexpr int N = 128; + std::array input; + std::array output; + std::iota(input.begin(), input.end(), 0); + std::fill(output.begin(), output.end(), 0); + + test(q, input, output, sycl::plus<>(), 0); + test(q, input, output, sycl::minimum<>(), + std::numeric_limits::max()); + test(q, input, output, sycl::maximum<>(), + std::numeric_limits::lowest()); + + test(q, input, output, sycl::plus(), 0); + test(q, input, output, sycl::minimum(), + std::numeric_limits::max()); + test(q, input, output, sycl::maximum(), + std::numeric_limits::lowest()); + +#ifdef SPIRV_1_3 + test(q, input, output, + sycl::multiplies(), 1); + test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_xor(), + 0); + test(q, input, output, sycl::bit_and(), ~0); +#endif // SPIRV_1_3 + + std::cout << "Test passed." << std::endl; +} diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp new file mode 100644 index 0000000000..4ad407cf5c --- /dev/null +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -0,0 +1,98 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3. +// That requires either adding a switch to clang (-spirv-max-version=1.3) or +// raising the spirv version from 1.1. to 1.3 for spirv translator +// unconditionally. Using operators specific for spirv 1.3 and higher with +// -spirv-max-version=1.1 being set by default causes assert/check fails +// in spirv translator. +// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ + %t13.out + +#include "support.h" +#include +#include +#include +#include +#include +using namespace sycl; + +template +void test(queue q, InputContainer input, OutputContainer output, + BinaryOperation binary_op, + typename OutputContainer::value_type identity) { + typedef typename InputContainer::value_type InputT; + typedef typename OutputContainer::value_type OutputT; + OutputT init = 42; + size_t N = input.size(); + size_t G = 64; + { + buffer in_buf(input.data(), input.size()); + buffer out_buf(output.data(), output.size()); + + q.submit([&](handler &cgh) { + accessor in{in_buf, cgh, sycl::read_only}; + accessor out{out_buf, cgh, sycl::write_only, sycl::no_init}; + cgh.parallel_for( + nd_range<1>(G, G), [=](nd_item<1> it) { + group<1> g = it.get_group(); + int lid = it.get_local_id(0); + out[0] = reduce_over_group(g, in[lid], binary_op); + out[1] = reduce_over_group(g, in[lid], init, binary_op); + out[2] = joint_reduce(g, in.get_pointer(), in.get_pointer() + N, + binary_op); + out[3] = joint_reduce(g, in.get_pointer(), in.get_pointer() + N, + init, binary_op); + }); + }); + } + // std::reduce is not implemented yet, so use std::accumulate instead + assert(output[0] == std::accumulate(input.begin(), input.begin() + G, + identity, binary_op)); + assert(output[1] == + std::accumulate(input.begin(), input.begin() + G, init, binary_op)); + assert(output[2] == + std::accumulate(input.begin(), input.end(), identity, binary_op)); + assert(output[3] == + std::accumulate(input.begin(), input.end(), init, binary_op)); +} + +int main() { + queue q; + if (!isSupportedDevice(q.get_device())) { + std::cout << "Skipping test\n"; + return 0; + } + + constexpr int N = 128; + std::array input; + std::array output; + std::iota(input.begin(), input.end(), 0); + std::fill(output.begin(), output.end(), 0); + + test(q, input, output, sycl::plus<>(), 0); + test(q, input, output, sycl::minimum<>(), + std::numeric_limits::max()); + test(q, input, output, sycl::maximum<>(), + std::numeric_limits::lowest()); + + test(q, input, output, sycl::plus(), 0); + test(q, input, output, sycl::minimum(), + std::numeric_limits::max()); + test(q, input, output, sycl::maximum(), + std::numeric_limits::lowest()); + +#ifdef SPIRV_1_3 + test(q, input, output, + sycl::multiplies(), 1); + test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), ~0); +#endif // SPIRV_1_3 + + std::cout << "Test passed." << std::endl; +} From 184ca52e6c75c5b85ba375a7944e5767a30391a3 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Fri, 4 Jun 2021 13:28:49 +0300 Subject: [PATCH 03/16] Clang-format fix --- SYCL/Basic/half_type.cpp | 37 ------------------- .../exclusive_scan_sycl2020.cpp | 10 +++-- .../inclusive_scan_sycl2020.cpp | 10 +++-- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 6 ++- 4 files changed, 16 insertions(+), 47 deletions(-) diff --git a/SYCL/Basic/half_type.cpp b/SYCL/Basic/half_type.cpp index 7465d45e3c..464f9ba627 100644 --- a/SYCL/Basic/half_type.cpp +++ b/SYCL/Basic/half_type.cpp @@ -192,44 +192,7 @@ inline bool bitwise_comparison_fp32(const half val, const uint32_t exp) { return reinterpret_cast(fp32) == exp; } -constexpr void constexpr_verify_add() { - constexpr half a{5.0}, b{2.0}, ref{7.0}; - constexpr half result = a + b; - constexpr half diff = result - ref; - static_assert(std::fabs(static_cast(diff)) < - std::numeric_limits::epsilon(), "Constexpr add is wrong"); -} - -constexpr void constexpr_verify_sub() { - constexpr half a{5.0f}, b{2.0}, ref{3.0}; - constexpr half result = a - b; - constexpr half diff = result - ref; - static_assert(std::fabs(static_cast(diff)) < - std::numeric_limits::epsilon(), "Constexpr sub is wrong"); -} - -constexpr void constexpr_verify_mul() { - constexpr half a{5.0f}, b{2.0}, ref{10.0}; - constexpr half result = a * b; - constexpr half diff = result - ref; - static_assert(std::fabs(static_cast(diff)) < - std::numeric_limits::epsilon(), "Constexpr mul is wrong"); -} - -constexpr void constexpr_verify_div() { - constexpr half a{5.0f}, b{2.0}, ref{2.5}; - constexpr half result = a / b; - constexpr half diff = result - ref; - static_assert(std::fabs(static_cast(diff)) < - std::numeric_limits::epsilon(), "Constexpr div is wrong"); -} - int main() { - constexpr_verify_add(); - constexpr_verify_sub(); - constexpr_verify_mul(); - constexpr_verify_div(); - device dev{default_selector()}; if (!dev.is_host() && !dev.has_extension("cl_khr_fp16")) { std::cout << "This device doesn't support the extension cl_khr_fp16" diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 1987e52656..3b5f28a9ce 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -151,11 +151,13 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, sycl::multiplies(), - 1); + test(q, input, output, + sycl::multiplies(), 1); test(q, input, output, sycl::bit_or(), 0); - test(q, input, output, sycl::bit_xor(), 0); - test(q, input, output, sycl::bit_and(), ~0); + test(q, input, output, sycl::bit_xor(), + 0); + test(q, input, output, sycl::bit_and(), + ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index f855c0717e..62121f6597 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -153,10 +153,12 @@ int main() { #ifdef SPIRV_1_3 test(q, input, output, sycl::multiplies(), 1); - test(q, input, output, sycl::bit_or(), 0); - test(q, input, output, sycl::bit_xor(), - 0); - test(q, input, output, sycl::bit_and(), ~0); + test(q, input, output, + sycl::bit_or(), 0); + test(q, input, output, + sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), + ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index 4ad407cf5c..53450b1b51 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -89,9 +89,11 @@ int main() { #ifdef SPIRV_1_3 test(q, input, output, sycl::multiplies(), 1); - test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_or(), + 0); test(q, input, output, sycl::bit_xor(), 0); - test(q, input, output, sycl::bit_and(), ~0); + test(q, input, output, + sycl::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; From 616cf45f0bd5679735bdc333a62c4844dad8696c Mon Sep 17 00:00:00 2001 From: mdimakov Date: Mon, 7 Jun 2021 11:03:28 +0300 Subject: [PATCH 04/16] Added std::inclusive/exclusive_scan --- .../exclusive_scan_sycl2020.cpp | 25 +++---------------- .../inclusive_scan_sycl2020.cpp | 25 +++---------------- 2 files changed, 8 insertions(+), 42 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 3b5f28a9ce..9edb0801c8 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -24,23 +24,6 @@ using namespace sycl; template class exclusive_scan_kernel; -// std::exclusive_scan isn't implemented yet, so use serial implementation -// instead -namespace emu { -template -OutputIterator exclusive_scan(InputIterator first, InputIterator last, - OutputIterator result, T init, - BinaryOperation binary_op) { - T partial = init; - for (InputIterator it = first; it != last; ++it) { - *(result++) = partial; - partial = binary_op(partial, *it); - } - return result; -} -} // namespace emu - template void test(queue q, InputContainer input, OutputContainer output, @@ -69,7 +52,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), + std::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), identity, binary_op); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -86,7 +69,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init, + std::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init, binary_op); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -103,7 +86,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), + std::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), identity, binary_op); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); @@ -120,7 +103,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init, + std::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init, binary_op); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); } diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 62121f6597..dd3e3eed81 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -24,23 +24,6 @@ using namespace sycl; template class inclusive_scan_kernel; -// std::inclusive_scan isn't implemented yet, so use serial implementation -// instead -namespace emu { -template -OutputIterator inclusive_scan(InputIterator first, InputIterator last, - OutputIterator result, BinaryOperation binary_op, - T init) { - T partial = init; - for (InputIterator it = first; it != last; ++it) { - partial = binary_op(partial, *it); - *(result++) = partial; - } - return result; -} -} // namespace emu - template void test(queue q, InputContainer input, OutputContainer output, @@ -69,7 +52,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), + std::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), binary_op, identity); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -86,7 +69,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), + std::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), binary_op, init); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -103,7 +86,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), + std::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), binary_op, identity); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); @@ -120,7 +103,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), + std::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), binary_op, init); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); } From 3835a718a5a68a92cc22846e518ada4093d3af5b Mon Sep 17 00:00:00 2001 From: mdimakov Date: Mon, 7 Jun 2021 16:08:31 +0300 Subject: [PATCH 05/16] Addressed review comments --- .../exclusive_scan_sycl2020.cpp | 30 +++++++++++++++---- .../inclusive_scan_sycl2020.cpp | 28 ++++++++++++++--- 2 files changed, 49 insertions(+), 9 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 9edb0801c8..de8519e633 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -32,13 +32,13 @@ void test(queue q, InputContainer input, OutputContainer output, typedef typename InputContainer::value_type InputT; typedef typename OutputContainer::value_type OutputT; typedef class exclusive_scan_kernel kernel_name0; - typedef class exclusive_scan_kernel kernel_name1; - typedef class exclusive_scan_kernel kernel_name2; - typedef class exclusive_scan_kernel kernel_name3; - OutputT init = 42; - size_t N = input.size(); size_t G = 64; + size_t N = input.size(); std::vector expected(N); + + // checking + // template + // T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -56,6 +56,12 @@ void test(queue q, InputContainer input, OutputContainer output, identity, binary_op); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + typedef class exclusive_scan_kernel kernel_name1; + OutputT init = 42; + + // checking + // template + // T exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -73,6 +79,13 @@ void test(queue q, InputContainer input, OutputContainer output, binary_op); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + typedef class exclusive_scan_kernel kernel_name2; + + // checking + // template + // OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + // BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -90,6 +103,13 @@ void test(queue q, InputContainer input, OutputContainer output, identity, binary_op); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); + typedef class exclusive_scan_kernel kernel_name3; + + // checking + // template + // OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, + // BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index dd3e3eed81..42bfb166a4 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -32,13 +32,13 @@ void test(queue q, InputContainer input, OutputContainer output, typedef typename InputContainer::value_type InputT; typedef typename OutputContainer::value_type OutputT; typedef class inclusive_scan_kernel kernel_name0; - typedef class inclusive_scan_kernel kernel_name1; - typedef class inclusive_scan_kernel kernel_name2; - typedef class inclusive_scan_kernel kernel_name3; - OutputT init = 42; size_t N = input.size(); size_t G = 64; std::vector expected(N); + + // checking + // template + // T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -56,6 +56,12 @@ void test(queue q, InputContainer input, OutputContainer output, binary_op, identity); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + typedef class inclusive_scan_kernel kernel_name1; + OutputT init = 42; + + // checking + // template + // T inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -73,6 +79,13 @@ void test(queue q, InputContainer input, OutputContainer output, binary_op, init); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); + typedef class inclusive_scan_kernel kernel_name2; + + // checking + // template + // OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + // BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -90,6 +103,13 @@ void test(queue q, InputContainer input, OutputContainer output, binary_op, identity); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); + typedef class inclusive_scan_kernel kernel_name3; + + // checking + // template + // OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + // BinaryOperation binary_op, T init) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); From 3ef82a3b0f2554ec74ab61dc1a3c12f2322f3b35 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Mon, 7 Jun 2021 16:16:15 +0300 Subject: [PATCH 06/16] Clang-format fix --- SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp | 17 ++++++++++------- SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp | 17 ++++++++++------- 2 files changed, 20 insertions(+), 14 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index de8519e633..8bb4fdda84 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -36,7 +36,7 @@ void test(queue q, InputContainer input, OutputContainer output, size_t N = input.size(); std::vector expected(N); - // checking + // checking // template // T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { @@ -59,9 +59,10 @@ void test(queue q, InputContainer input, OutputContainer output, typedef class exclusive_scan_kernel kernel_name1; OutputT init = 42; - // checking + // checking // template - // T exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) + // T exclusive_scan_over_group(Group g, V x, T init, BinaryOperation + // binary_op) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -81,10 +82,11 @@ void test(queue q, InputContainer input, OutputContainer output, typedef class exclusive_scan_kernel kernel_name2; - // checking + // checking // template - // OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + // OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr + // result, // BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); @@ -105,10 +107,11 @@ void test(queue q, InputContainer input, OutputContainer output, typedef class exclusive_scan_kernel kernel_name3; - // checking + // checking // template - // OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, + // OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr + // result, T init, // BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 42bfb166a4..3ecdefc559 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -36,7 +36,7 @@ void test(queue q, InputContainer input, OutputContainer output, size_t G = 64; std::vector expected(N); - // checking + // checking // template // T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { @@ -59,9 +59,10 @@ void test(queue q, InputContainer input, OutputContainer output, typedef class inclusive_scan_kernel kernel_name1; OutputT init = 42; - // checking + // checking // template - // T inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) + // T inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T + // init) { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -81,10 +82,11 @@ void test(queue q, InputContainer input, OutputContainer output, typedef class inclusive_scan_kernel kernel_name2; - // checking + // checking // template - // OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + // OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr + // result, // BinaryOperation binary_op) { buffer in_buf(input.data(), input.size()); @@ -105,10 +107,11 @@ void test(queue q, InputContainer input, OutputContainer output, typedef class inclusive_scan_kernel kernel_name3; - // checking + // checking // template - // OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + // OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr + // result, // BinaryOperation binary_op, T init) { buffer in_buf(input.data(), input.size()); From 04ba6c0ab88b6c135c962ce6c00fcd109359db2f Mon Sep 17 00:00:00 2001 From: mdimakov Date: Tue, 8 Jun 2021 18:59:20 +0300 Subject: [PATCH 07/16] Renamed kersome kernels --- SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp | 14 +++++++------- SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp | 16 ++++++++-------- 2 files changed, 15 insertions(+), 15 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 8bb4fdda84..a39098ab98 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -157,13 +157,13 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - sycl::multiplies(), 1); - test(q, input, output, sycl::bit_or(), 0); - test(q, input, output, sycl::bit_xor(), - 0); - test(q, input, output, sycl::bit_and(), - ~0); + test(q, input, output, + sycl::multiplies(), 1); + test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_xor(), + 0); + test(q, input, output, sycl::bit_and(), + ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 3ecdefc559..e91148e5ff 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -157,14 +157,14 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - sycl::multiplies(), 1); - test(q, input, output, - sycl::bit_or(), 0); - test(q, input, output, - sycl::bit_xor(), 0); - test(q, input, output, sycl::bit_and(), - ~0); + test(q, input, output, + sycl::multiplies(), 1); + test(q, input, output, + sycl::bit_or(), 0); + test(q, input, output, + sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), + ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; From 9e9c2eecffe06d700b16568a2d98b8210bdd2f68 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Wed, 9 Jun 2021 09:46:01 +0300 Subject: [PATCH 08/16] Added constexpr to variables and renamed kernel names --- .../exclusive_scan_sycl2020.cpp | 6 +++--- .../inclusive_scan_sycl2020.cpp | 6 +++--- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 20 +++++++++---------- 3 files changed, 16 insertions(+), 16 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index a39098ab98..dc7eb48fda 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -32,8 +32,8 @@ void test(queue q, InputContainer input, OutputContainer output, typedef typename InputContainer::value_type InputT; typedef typename OutputContainer::value_type OutputT; typedef class exclusive_scan_kernel kernel_name0; - size_t G = 64; - size_t N = input.size(); + constexpr size_t G = 64; + constexpr size_t N = input.size(); std::vector expected(N); // checking @@ -57,7 +57,7 @@ void test(queue q, InputContainer input, OutputContainer output, assert(std::equal(output.begin(), output.begin() + G, expected.begin())); typedef class exclusive_scan_kernel kernel_name1; - OutputT init = 42; + constexpr OutputT init = 42; // checking // template diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index e91148e5ff..2810cf7326 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -32,8 +32,8 @@ void test(queue q, InputContainer input, OutputContainer output, typedef typename InputContainer::value_type InputT; typedef typename OutputContainer::value_type OutputT; typedef class inclusive_scan_kernel kernel_name0; - size_t N = input.size(); - size_t G = 64; + constexpr size_t N = input.size(); + constexpr size_t G = 64; std::vector expected(N); // checking @@ -57,7 +57,7 @@ void test(queue q, InputContainer input, OutputContainer output, assert(std::equal(output.begin(), output.begin() + G, expected.begin())); typedef class inclusive_scan_kernel kernel_name1; - OutputT init = 42; + constexpr OutputT init = 42; // checking // template diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index 53450b1b51..98b6ff59ae 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -27,9 +27,9 @@ void test(queue q, InputContainer input, OutputContainer output, typename OutputContainer::value_type identity) { typedef typename InputContainer::value_type InputT; typedef typename OutputContainer::value_type OutputT; - OutputT init = 42; - size_t N = input.size(); - size_t G = 64; + constexpr OutputT init = 42; + constexpr size_t N = input.size(); + constexpr size_t G = 64; { buffer in_buf(input.data(), input.size()); buffer out_buf(output.data(), output.size()); @@ -87,13 +87,13 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - sycl::multiplies(), 1); - test(q, input, output, sycl::bit_or(), - 0); - test(q, input, output, sycl::bit_xor(), 0); - test(q, input, output, - sycl::bit_and(), ~0); + test(q, input, output, + sycl::multiplies(), 1); + test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_xor(), + 0); + test(q, input, output, sycl::bit_and(), + ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; From c8f5c9d9e04ec197369483fb62a131d51b254aa8 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Wed, 9 Jun 2021 09:48:37 +0300 Subject: [PATCH 09/16] Clang-format fix --- SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp | 10 ++++------ SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp | 13 +++++-------- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 10 ++++------ 3 files changed, 13 insertions(+), 20 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index dc7eb48fda..bf1b680bf3 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -157,13 +157,11 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - sycl::multiplies(), 1); + test(q, input, output, sycl::multiplies(), + 1); test(q, input, output, sycl::bit_or(), 0); - test(q, input, output, sycl::bit_xor(), - 0); - test(q, input, output, sycl::bit_and(), - ~0); + test(q, input, output, sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 2810cf7326..164d502d12 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -157,14 +157,11 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - sycl::multiplies(), 1); - test(q, input, output, - sycl::bit_or(), 0); - test(q, input, output, - sycl::bit_xor(), 0); - test(q, input, output, sycl::bit_and(), - ~0); + test(q, input, output, sycl::multiplies(), + 1); + test(q, input, output, sycl::bit_or(), 0); + test(q, input, output, sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index 98b6ff59ae..dcc2cd984a 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -87,13 +87,11 @@ int main() { std::numeric_limits::lowest()); #ifdef SPIRV_1_3 - test(q, input, output, - sycl::multiplies(), 1); + test(q, input, output, sycl::multiplies(), + 1); test(q, input, output, sycl::bit_or(), 0); - test(q, input, output, sycl::bit_xor(), - 0); - test(q, input, output, sycl::bit_and(), - ~0); + test(q, input, output, sycl::bit_xor(), 0); + test(q, input, output, sycl::bit_and(), ~0); #endif // SPIRV_1_3 std::cout << "Test passed." << std::endl; From d02f4e0a59b93166eddcecb71959c3560cf5c57f Mon Sep 17 00:00:00 2001 From: mdimakov Date: Wed, 9 Jun 2021 15:04:56 +0300 Subject: [PATCH 10/16] Use std::reduce instead of std::accumulate --- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index dcc2cd984a..d8e619fd5f 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -50,15 +50,14 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - // std::reduce is not implemented yet, so use std::accumulate instead - assert(output[0] == std::accumulate(input.begin(), input.begin() + G, + assert(output[0] == std::reduce(input.begin(), input.begin() + G, identity, binary_op)); assert(output[1] == - std::accumulate(input.begin(), input.begin() + G, init, binary_op)); + std::reduce(input.begin(), input.begin() + G, init, binary_op)); assert(output[2] == - std::accumulate(input.begin(), input.end(), identity, binary_op)); + std::reduce(input.begin(), input.end(), identity, binary_op)); assert(output[3] == - std::accumulate(input.begin(), input.end(), init, binary_op)); + std::reduce(input.begin(), input.end(), init, binary_op)); } int main() { From 50ed2a1f110bac255c4a274e78174b6f8f85441d Mon Sep 17 00:00:00 2001 From: mdimakov Date: Tue, 15 Jun 2021 11:20:32 +0300 Subject: [PATCH 11/16] Clang-format fix --- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index d8e619fd5f..a54ece3e39 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -50,14 +50,13 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - assert(output[0] == std::reduce(input.begin(), input.begin() + G, - identity, binary_op)); + assert(output[0] == + std::reduce(input.begin(), input.begin() + G, identity, binary_op)); assert(output[1] == std::reduce(input.begin(), input.begin() + G, init, binary_op)); assert(output[2] == std::reduce(input.begin(), input.end(), identity, binary_op)); - assert(output[3] == - std::reduce(input.begin(), input.end(), init, binary_op)); + assert(output[3] == std::reduce(input.begin(), input.end(), init, binary_op)); } int main() { From db121968563eb190d8282904670e8d3e1f15553c Mon Sep 17 00:00:00 2001 From: mdimakov Date: Thu, 17 Jun 2021 09:30:26 +0300 Subject: [PATCH 12/16] Replaced std functions --- .../exclusive_scan_sycl2020.cpp | 25 ++++++++++++++++--- .../inclusive_scan_sycl2020.cpp | 25 ++++++++++++++++--- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 5 ++-- 3 files changed, 45 insertions(+), 10 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index bf1b680bf3..9cec348e5f 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -24,6 +24,23 @@ using namespace sycl; template class exclusive_scan_kernel; +// std::exclusive_scan isn't implemented yet, so use serial implementation +// instead +namespace emu { +template +OutputIterator exclusive_scan(InputIterator first, InputIterator last, + OutputIterator result, T init, + BinaryOperation binary_op) { + T partial = init; + for (InputIterator it = first; it != last; ++it) { + *(result++) = partial; + partial = binary_op(partial, *it); + } + return result; +} +} // namespace emu + template void test(queue q, InputContainer input, OutputContainer output, @@ -52,7 +69,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), + emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), identity, binary_op); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -76,7 +93,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init, + emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init, binary_op); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -101,7 +118,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), + emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), identity, binary_op); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); @@ -126,7 +143,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init, + emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init, binary_op); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); } diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 164d502d12..b03db6124d 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -24,6 +24,23 @@ using namespace sycl; template class inclusive_scan_kernel; +// std::inclusive_scan isn't implemented yet, so use serial implementation +// instead +namespace emu { +template +OutputIterator inclusive_scan(InputIterator first, InputIterator last, + OutputIterator result, BinaryOperation binary_op, + T init) { + T partial = init; + for (InputIterator it = first; it != last; ++it) { + partial = binary_op(partial, *it); + *(result++) = partial; + } + return result; +} +} // namespace emu + template void test(queue q, InputContainer input, OutputContainer output, @@ -52,7 +69,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), + emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), binary_op, identity); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -76,7 +93,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), + emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(), binary_op, init); assert(std::equal(output.begin(), output.begin() + G, expected.begin())); @@ -101,7 +118,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), + emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), binary_op, identity); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); @@ -126,7 +143,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - std::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), + emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(), binary_op, init); assert(std::equal(output.begin(), output.begin() + N, expected.begin())); } diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index a54ece3e39..cd58b878b5 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -50,8 +50,9 @@ void test(queue q, InputContainer input, OutputContainer output, }); }); } - assert(output[0] == - std::reduce(input.begin(), input.begin() + G, identity, binary_op)); + // std::reduce is not implemented yet, so use std::accumulate instead + assert(output[0] == std::accumulate(input.begin(), input.begin() + G, + identity, binary_op)); assert(output[1] == std::reduce(input.begin(), input.begin() + G, init, binary_op)); assert(output[2] == From 719eb3e54b786dee7cc44fdb59c6a975d36ef4ae Mon Sep 17 00:00:00 2001 From: mdimakov Date: Thu, 17 Jun 2021 10:25:02 +0300 Subject: [PATCH 13/16] Removed the using of std::reduce --- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index cd58b878b5..dcc2cd984a 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -54,10 +54,11 @@ void test(queue q, InputContainer input, OutputContainer output, assert(output[0] == std::accumulate(input.begin(), input.begin() + G, identity, binary_op)); assert(output[1] == - std::reduce(input.begin(), input.begin() + G, init, binary_op)); + std::accumulate(input.begin(), input.begin() + G, init, binary_op)); assert(output[2] == - std::reduce(input.begin(), input.end(), identity, binary_op)); - assert(output[3] == std::reduce(input.begin(), input.end(), init, binary_op)); + std::accumulate(input.begin(), input.end(), identity, binary_op)); + assert(output[3] == + std::accumulate(input.begin(), input.end(), init, binary_op)); } int main() { From bfaaa0373981a7958390cf349ed4c1e43d95cb95 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Fri, 18 Jun 2021 09:28:56 +0300 Subject: [PATCH 14/16] Addressed review comments --- SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp | 9 ++------- SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp | 9 ++------- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 9 ++------- 3 files changed, 6 insertions(+), 21 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 9cec348e5f..18537c3b2b 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -3,13 +3,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3. -// That requires either adding a switch to clang (-spirv-max-version=1.3) or -// raising the spirv version from 1.1. to 1.3 for spirv translator -// unconditionally. Using operators specific for spirv 1.3 and higher with -// -spirv-max-version=1.1 being set by default causes assert/check fails -// in spirv translator. -// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ %t13.out #include "support.h" @@ -26,6 +20,7 @@ class exclusive_scan_kernel; // std::exclusive_scan isn't implemented yet, so use serial implementation // instead +// TODO: use std::exclusive_scan when it will be supported namespace emu { template diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index b03db6124d..2cd0cc4fba 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -3,13 +3,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3. -// That requires either adding a switch to clang (-spirv-max-version=1.3) or -// raising the spirv version from 1.1. to 1.3 for spirv translator -// unconditionally. Using operators specific for spirv 1.3 and higher with -// -spirv-max-version=1.1 being set by default causes assert/check fails -// in spirv translator. -// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ %t13.out #include "support.h" @@ -26,6 +20,7 @@ class inclusive_scan_kernel; // std::inclusive_scan isn't implemented yet, so use serial implementation // instead +// TODO: use std::inclusive_scan when it will be supported namespace emu { template diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index dcc2cd984a..bf24bfff4e 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -3,13 +3,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3. -// That requires either adding a switch to clang (-spirv-max-version=1.3) or -// raising the spirv version from 1.1. to 1.3 for spirv translator -// unconditionally. Using operators specific for spirv 1.3 and higher with -// -spirv-max-version=1.1 being set by default causes assert/check fails -// in spirv translator. -// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ %t13.out #include "support.h" @@ -51,6 +45,7 @@ void test(queue q, InputContainer input, OutputContainer output, }); } // std::reduce is not implemented yet, so use std::accumulate instead + // TODO: use std::reduce when it will be supported assert(output[0] == std::accumulate(input.begin(), input.begin() + G, identity, binary_op)); assert(output[1] == From e4dd6bfc1b73da19aa70808916053c86160648fd Mon Sep 17 00:00:00 2001 From: mdimakov Date: Fri, 18 Jun 2021 12:14:31 +0300 Subject: [PATCH 15/16] Fix RUN lines --- SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp | 3 +-- SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp | 3 +-- SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 3 +-- 3 files changed, 3 insertions(+), 6 deletions(-) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 18537c3b2b..3a6ad67e95 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -3,8 +3,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ - %t13.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out #include "support.h" #include diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 2cd0cc4fba..d7b2327ef7 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -3,8 +3,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ - %t13.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out #include "support.h" #include diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index bf24bfff4e..ca63dbad53 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -3,8 +3,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \ - %t13.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out #include "support.h" #include From 9c756051777ffd27a0a8aebb910a4a08a6e58c46 Mon Sep 17 00:00:00 2001 From: mdimakov Date: Fri, 18 Jun 2021 17:30:46 +0300 Subject: [PATCH 16/16] Unsupported cuda --- SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp | 1 + SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp | 1 + SYCL/GroupAlgorithm/reduce_sycl2020.cpp | 1 + 3 files changed, 3 insertions(+) diff --git a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 3a6ad67e95..5424cb9d31 100644 --- a/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -1,3 +1,4 @@ +// UNSUPPORTED: cuda // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp index d7b2327ef7..c9c76e06fb 100644 --- a/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -1,3 +1,4 @@ +// UNSUPPORTED: cuda // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp index ca63dbad53..82259c7435 100644 --- a/SYCL/GroupAlgorithm/reduce_sycl2020.cpp +++ b/SYCL/GroupAlgorithm/reduce_sycl2020.cpp @@ -1,3 +1,4 @@ +// UNSUPPORTED: cuda // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out