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

Commit 2db68c1

Browse files
cperkinsintelcperkinsintel
authored
[SYCL] test new complex support by group_algorithms (#767)
Signed-off-by: Chris Perkins <[email protected]>
1 parent 5ecf10d commit 2db68c1

File tree

3 files changed

+96
-6
lines changed

3 files changed

+96
-6
lines changed

‎SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp

Lines changed: 32 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
11
// UNSUPPORTED: cuda || hip
2-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
33
// RUN: %CPU_RUN_PLACEHOLDER %t.out
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66

7-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out
7+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
88

99
#include "support.h"
1010
#include <CL/sycl.hpp>
@@ -186,5 +186,35 @@ int main() {
186186
sycl::bit_and<int>(), ~0);
187187
#endif // SPIRV_1_3
188188

189+
// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
190+
// https://github.com/intel/llvm/pull/5108/ ) joint_exclusive_scan and
191+
// exclusive_scan_over_group now operate on std::complex but limited to the
192+
// sycl::plus binary operation.
193+
#ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
194+
std::array<std::complex<float>, N> input_cf;
195+
std::array<std::complex<float>, N> output_cf;
196+
std::iota(input_cf.begin(), input_cf.end(), 0);
197+
std::fill(output_cf.begin(), output_cf.end(), 0);
198+
test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
199+
sycl::plus<std::complex<float>>(), 0);
200+
test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);
201+
202+
if (q.get_device().has(aspect::fp64)) {
203+
std::array<std::complex<double>, N> input_cd;
204+
std::array<std::complex<double>, N> output_cd;
205+
std::iota(input_cd.begin(), input_cd.end(), 0);
206+
std::fill(output_cd.begin(), output_cd.end(), 0);
207+
test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
208+
sycl::plus<std::complex<double>>(), 0);
209+
test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
210+
0);
211+
} else {
212+
std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
213+
<< std::endl;
214+
}
215+
#else
216+
static_assert(false, "SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS not defined");
217+
#endif
218+
189219
std::cout << "Test passed." << std::endl;
190220
}

‎SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp

Lines changed: 31 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
11
// UNSUPPORTED: cuda || hip
2-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
33
// RUN: %CPU_RUN_PLACEHOLDER %t.out
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66

7-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out
7+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
88

99
#include "support.h"
1010
#include <CL/sycl.hpp>
@@ -186,5 +186,34 @@ int main() {
186186
sycl::bit_and<int>(), ~0);
187187
#endif // SPIRV_1_3
188188

189+
// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
190+
// https://github.com/intel/llvm/pull/5108/ ) joint_inclusive_scan and
191+
// inclusive_scan_over_group now operate on std::complex limited to using the
192+
// sycl::plus binary operation.
193+
#ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
194+
std::array<std::complex<float>, N> input_cf;
195+
std::array<std::complex<float>, N> output_cf;
196+
std::iota(input_cf.begin(), input_cf.end(), 0);
197+
std::fill(output_cf.begin(), output_cf.end(), 0);
198+
test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
199+
sycl::plus<std::complex<float>>(), 0);
200+
test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);
201+
202+
if (q.get_device().has(aspect::fp64)) {
203+
std::array<std::complex<double>, N> input_cd;
204+
std::array<std::complex<double>, N> output_cd;
205+
std::iota(input_cd.begin(), input_cd.end(), 0);
206+
std::fill(output_cd.begin(), output_cd.end(), 0);
207+
test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
208+
sycl::plus<std::complex<double>>(), 0);
209+
test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
210+
0);
211+
} else {
212+
std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
213+
<< std::endl;
214+
}
215+
#else
216+
static_assert(false, "SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS not defined");
217+
#endif
189218
std::cout << "Test passed." << std::endl;
190219
}

‎SYCL/GroupAlgorithm/reduce_sycl2020.cpp

Lines changed: 33 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,16 @@
11
// UNSUPPORTED: cuda || hip
2-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %t.out
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
33
// RUN: %CPU_RUN_PLACEHOLDER %t.out
44
// RUN: %GPU_RUN_PLACEHOLDER %t.out
55
// RUN: %ACC_RUN_PLACEHOLDER %t.out
66

7-
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out
7+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
88

99
#include "support.h"
1010
#include <CL/sycl.hpp>
1111
#include <algorithm>
1212
#include <cassert>
13+
#include <complex>
1314
#include <limits>
1415
#include <numeric>
1516
using namespace sycl;
@@ -89,5 +90,35 @@ int main() {
8990
test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
9091
#endif // SPIRV_1_3
9192

93+
// as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
94+
// https://github.com/intel/llvm/pull/5108/ ) joint_reduce and
95+
// reduce_over_group now operate on std::complex limited to using the
96+
// sycl::plus binary operation.
97+
#ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
98+
std::array<std::complex<float>, N> input_cf;
99+
std::array<std::complex<float>, 4> output_cf;
100+
std::iota(input_cf.begin(), input_cf.end(), 0);
101+
std::fill(output_cf.begin(), output_cf.end(), 0);
102+
test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
103+
sycl::plus<std::complex<float>>(), 0);
104+
test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);
105+
106+
if (q.get_device().has(aspect::fp64)) {
107+
std::array<std::complex<double>, N> input_cd;
108+
std::array<std::complex<double>, 4> output_cd;
109+
std::iota(input_cd.begin(), input_cd.end(), 0);
110+
std::fill(output_cd.begin(), output_cd.end(), 0);
111+
test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
112+
sycl::plus<std::complex<double>>(), 0);
113+
test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
114+
0);
115+
} else {
116+
std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
117+
<< std::endl;
118+
}
119+
#else
120+
static_assert(false, "SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS not defined");
121+
#endif
122+
92123
std::cout << "Test passed." << std::endl;
93124
}

0 commit comments

Comments
 (0)