|
| 1 | +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out |
| 2 | +// RUN: %CPU_RUN_PLACEHOLDER %t.out |
| 3 | +// RUN: %GPU_RUN_PLACEHOLDER %t.out |
| 4 | +// RUN: %ACC_RUN_PLACEHOLDER %t.out |
| 5 | + |
| 6 | +// This test checks that if the custom type supports operations like +=, then |
| 7 | +// such operations can be used for the reduction objects in kernels.. |
| 8 | + |
| 9 | +#include <CL/sycl.hpp> |
| 10 | +#include <iostream> |
| 11 | +#include <random> |
| 12 | + |
| 13 | +using namespace sycl; |
| 14 | +using namespace sycl::ONEAPI; |
| 15 | + |
| 16 | +struct XY { |
| 17 | + XY() : X(0), Y(0) {} |
| 18 | + XY(float X, float Y) : X(X), Y(Y) {} |
| 19 | + float X; |
| 20 | + float Y; |
| 21 | + float x() const { return X; }; |
| 22 | + float y() const { return Y; }; |
| 23 | + XY &operator+=(const XY &RHS) { |
| 24 | + X += RHS.X; |
| 25 | + Y += RHS.Y; |
| 26 | + return *this; |
| 27 | + } |
| 28 | +}; |
| 29 | + |
| 30 | +void setXY(float2 &Data, float X, float Y) { Data = {X, Y}; } |
| 31 | +void setXY(XY &Data, float X, float Y) { |
| 32 | + Data.X = X; |
| 33 | + Data.Y = Y; |
| 34 | +} |
| 35 | + |
| 36 | +template <typename Name, typename T, typename BinaryOperation> |
| 37 | +int test(T Identity, BinaryOperation BOp) { |
| 38 | + constexpr size_t N = 16; |
| 39 | + constexpr size_t L = 4; |
| 40 | + |
| 41 | + queue Q; |
| 42 | + T *Data = malloc_shared<T>(N, Q); |
| 43 | + T *Res = malloc_shared<T>(1, Q); |
| 44 | + T Expected = Identity; |
| 45 | + for (size_t I = 0; I < N; I++) { |
| 46 | + setXY(Data[I], I, I + 1); |
| 47 | + setXY(Expected, Expected.x() + I, Expected.y() + I + 1); |
| 48 | + } |
| 49 | + |
| 50 | + *Res = Identity; |
| 51 | + auto Red = reduction(Res, Identity, BOp); |
| 52 | + Q.submit([&](handler &H) { |
| 53 | + H.parallel_for<Name>(nd_range<1>{N, L}, Red, |
| 54 | + [=](nd_item<1> ID, auto &Sum) { |
| 55 | + size_t GID = ID.get_global_id(0); |
| 56 | + Sum += Data[GID]; |
| 57 | + }); |
| 58 | + }).wait(); |
| 59 | + |
| 60 | + int Error = 0; |
| 61 | + if (Expected.x() != Res->x() || Expected.y() != Res->y()) { |
| 62 | + std::cerr << "Error: expected = (" << Expected.x() << ", " << Expected.y() |
| 63 | + << "); computed = (" << Res->x() << ", " << Res->y() << ")\n"; |
| 64 | + Error = 1; |
| 65 | + } |
| 66 | + free(Res, Q); |
| 67 | + free(Data, Q); |
| 68 | + return Error; |
| 69 | +} |
| 70 | + |
| 71 | +int main() { |
| 72 | + int Error = 0; |
| 73 | + Error += test<class A, float2>(float2{}, std::plus<>{}); |
| 74 | + Error += test<class B, XY>( |
| 75 | + XY{}, [](auto A, auto B) { return XY(A.X + B.X, A.Y + B.Y); }); |
| 76 | + if (!Error) |
| 77 | + std::cout << "Passed.\n"; |
| 78 | + return Error; |
| 79 | +} |
0 commit comments