Skip to content

Commit

Permalink
[SYCL] test new complex support by group_algorithms (intel/llvm-test-…
Browse files Browse the repository at this point in the history
…suite#767)

Signed-off-by: Chris Perkins <[email protected]>
  • Loading branch information
cperkinsintel authored Feb 10, 2022
1 parent 1a30328 commit df70e99
Show file tree
Hide file tree
Showing 3 changed files with 96 additions and 6 deletions.
34 changes: 32 additions & 2 deletions SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp
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
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
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;
}

0 comments on commit df70e99

Please sign in to comment.