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

[SYCL] test new complex support by group_algorithms #767

Merged
merged 8 commits into from
Feb 10, 2022
Merged

[SYCL] test new complex support by group_algorithms #767

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
d6f261d
updating tests to reflect new support of std::complex by joint_reduce…
cperkinsintel Jan 25, 2022
2d337e8
exclusive_scan tests expanded
cperkinsintel Jan 26, 2022
8885bde
inclusive_scan_over_group and joint_inclusive_scan
cperkinsintel Jan 26, 2022
dbe7f52
new testing for plus<>, double
cperkinsintel Feb 1, 2022
d07c067
check SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS and aspect::fp64
cperkinsintel Feb 1, 2022
6ac9008
device_code_split=per_kernel
cperkinsintel Feb 9, 2022
f6a36a8
Merge branch 'intel' into cperkins-complex-group-algo-support
cperkinsintel Feb 9, 2022
69277ff
re-run tests by changing comment
cperkinsintel Feb 9, 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
34 changes: 32 additions & 2 deletions SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,10 +1,10 @@
// UNSUPPORTED: cuda || hip
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// 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 -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out

#include "support.h"
#include <CL/sycl.hpp>
Expand Down Expand Up @@ -186,5 +186,35 @@ int main() {
sycl::bit_and<int>(), ~0);
#endif // SPIRV_1_3

// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
// https://github.com/intel/llvm/pull/5108/ ) joint_exclusive_scan and
// exclusive_scan_over_group now operate on std::complex but limited to the
// sycl::plus binary operation.
#ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
std::array<std::complex<float>, N> input_cf;
std::array<std::complex<float>, N> output_cf;
std::iota(input_cf.begin(), input_cf.end(), 0);
std::fill(output_cf.begin(), output_cf.end(), 0);
test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
sycl::plus<std::complex<float>>(), 0);
test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);

if (q.get_device().has(aspect::fp64)) {
std::array<std::complex<double>, N> input_cd;
std::array<std::complex<double>, N> output_cd;
std::iota(input_cd.begin(), input_cd.end(), 0);
std::fill(output_cd.begin(), output_cd.end(), 0);
test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
sycl::plus<std::complex<double>>(), 0);
test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
0);
} else {
std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
<< std::endl;
}
#else
static_assert(false, "SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS not defined");
#endif

std::cout << "Test passed." << std::endl;
}
33 changes: 31 additions & 2 deletions SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,10 +1,10 @@
// UNSUPPORTED: cuda || hip
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// 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 -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out

#include "support.h"
#include <CL/sycl.hpp>
Expand Down Expand Up @@ -186,5 +186,34 @@ int main() {
sycl::bit_and<int>(), ~0);
#endif // SPIRV_1_3

// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
// https://github.com/intel/llvm/pull/5108/ ) joint_inclusive_scan and
// inclusive_scan_over_group now operate on std::complex limited to using the
// sycl::plus binary operation.
#ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
std::array<std::complex<float>, N> input_cf;
std::array<std::complex<float>, N> output_cf;
std::iota(input_cf.begin(), input_cf.end(), 0);
std::fill(output_cf.begin(), output_cf.end(), 0);
test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
sycl::plus<std::complex<float>>(), 0);
test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);

if (q.get_device().has(aspect::fp64)) {
std::array<std::complex<double>, N> input_cd;
std::array<std::complex<double>, N> output_cd;
std::iota(input_cd.begin(), input_cd.end(), 0);
std::fill(output_cd.begin(), output_cd.end(), 0);
test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
sycl::plus<std::complex<double>>(), 0);
test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
0);
} else {
std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
<< std::endl;
}
#else
static_assert(false, "SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS not defined");
#endif
std::cout << "Test passed." << std::endl;
}
35 changes: 33 additions & 2 deletions SYCL/GroupAlgorithm/reduce_sycl2020.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,15 +1,16 @@
// UNSUPPORTED: cuda || hip
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// 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 -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out

#include "support.h"
#include <CL/sycl.hpp>
#include <algorithm>
#include <cassert>
#include <complex>
#include <limits>
#include <numeric>
using namespace sycl;
Expand Down Expand Up @@ -89,5 +90,35 @@ int main() {
test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
#endif // SPIRV_1_3

// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
// https://github.com/intel/llvm/pull/5108/ ) joint_reduce and
// reduce_over_group now operate on std::complex limited to using the
// sycl::plus binary operation.
#ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
std::array<std::complex<float>, N> input_cf;
std::array<std::complex<float>, 4> output_cf;
std::iota(input_cf.begin(), input_cf.end(), 0);
std::fill(output_cf.begin(), output_cf.end(), 0);
test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
sycl::plus<std::complex<float>>(), 0);
test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);

if (q.get_device().has(aspect::fp64)) {
std::array<std::complex<double>, N> input_cd;
std::array<std::complex<double>, 4> output_cd;
std::iota(input_cd.begin(), input_cd.end(), 0);
std::fill(output_cd.begin(), output_cd.end(), 0);
test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
sycl::plus<std::complex<double>>(), 0);
test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
0);
} else {
std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
<< std::endl;
}
#else
static_assert(false, "SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS not defined");
#endif

std::cout << "Test passed." << std::endl;
}