Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[SYCL] Introduce sycl complex marray specialization #8647

Closed
Show file tree
Hide file tree
Changes from 43 commits
Commits
Show all changes
53 commits
Select commit Hold shift + click to select a range
9eca5bc
[SYCL][MARRAY] introduce sycl complex marray specialization
jle-quel Mar 14, 2023
c5ba5af
[SYCL][MARRAY] add marray specialization's tests
jle-quel Mar 14, 2023
98eb062
[SYCL][Complex] fix typo in complex's tests
jle-quel Mar 14, 2023
9a7570c
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Mar 28, 2023
56a9ddd
Add helper functions for marray complex testing
jle-quel Mar 28, 2023
ab1d887
Fix warning messages
jle-quel Mar 28, 2023
cbec286
Add test cases for marray complex
jle-quel Mar 28, 2023
5a5f647
Add marray complex getter tests
jle-quel Mar 28, 2023
7db90aa
Add marray complex math tests
jle-quel Mar 28, 2023
0a2b793
Add marray complex operator tests
jle-quel Mar 28, 2023
f9e3e76
Address PR's feedbacks by modifying the marray's constructors
jle-quel May 3, 2023
97f81e9
format code with the latest clang-format
jle-quel May 3, 2023
1d6669c
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Jun 27, 2023
d5b7664
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Jun 28, 2023
95c0453
Update marray tests with the latests %{build} and %{run}
jle-quel Jun 29, 2023
46079c2
Move marray's helper traits to detail's namespace, so exp::marray doe…
jle-quel Jun 29, 2023
c40fb02
Apply clang-format
jle-quel Jun 29, 2023
6fd924d
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Jul 12, 2023
5bced0f
Encapsulate marray's helper into struct which is a friend of class ma…
jle-quel Jul 24, 2023
1fb5a4f
Update marray's contructor tests
jle-quel Jul 24, 2023
91c14e5
format code with the latest clang-format
jle-quel Jul 24, 2023
0b54a77
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Jul 24, 2023
18b4277
Replace ComplexDataT by value_type
jle-quel Jul 25, 2023
1627da7
Remove -fsycl-device-code-split=per_kernel from the tests
jle-quel Jul 25, 2023
884b12e
Enable std::enable_if via Return type isntead of type template parame…
jle-quel Jul 25, 2023
79b2711
Run clang-format
jle-quel Jul 25, 2023
e4a1dbd
Remove special test cases for windows
jle-quel Aug 7, 2023
d2df5d1
Fix marray<std::complex<>> casting to never be a marray<std::complex<…
jle-quel Aug 7, 2023
c890407
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Aug 7, 2023
8d086dd
Use a less conflicting name for defining the marray's macros
jle-quel Aug 8, 2023
fcb3116
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Aug 8, 2023
f194bce
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Aug 16, 2023
58d19d3
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Sep 12, 2023
b92cce6
Apply clang-format
jle-quel Sep 12, 2023
b155c1d
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Oct 18, 2023
da12998
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Oct 31, 2023
a179509
Split marray header file and tests into multiple files
jle-quel Nov 2, 2023
8799339
Apply clang-format
jle-quel Nov 2, 2023
f756d38
Remove the real and imag member functions
jle-quel Nov 2, 2023
051db38
Add new real and imag free functions
jle-quel Nov 2, 2023
80b41ba
Update tests
jle-quel Nov 2, 2023
4dd8cd7
Apply clang-format
jle-quel Nov 2, 2023
95258e1
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel Feb 8, 2024
4728073
Rename HelperFlattenMArrayArg to FlattenMArrayArgHelper and update na…
jle-quel Feb 13, 2024
fd2759f
Update FIXME comments
jle-quel Feb 13, 2024
152156c
Make marray_data_t a private type alias
jle-quel Feb 13, 2024
735b2b8
Apply review's and new specification's changes
jle-quel May 6, 2024
0fed583
Apply clang-format
jle-quel May 7, 2024
5bbb307
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel May 8, 2024
38e2f9c
Rewrite tests that fail when usm_shared_allocations is not supported
jle-quel May 10, 2024
a3b4b05
Apply clang-format manually
jle-quel May 10, 2024
8a58080
Refactor marray builtins
jle-quel May 20, 2024
f38241a
Merge branch 'sycl' of github.com:jle-quel/llvm into jle-quel/introdu…
jle-quel May 21, 2024
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
Original file line number Diff line number Diff line change
Expand Up @@ -12,5 +12,7 @@

#include "./detail/complex.hpp"
#include "./detail/complex_math.hpp"
#include "./detail/marray.hpp"
#include "./detail/marray_math.hpp"

#endif // SYCL_EXT_ONEAPI_COMPLEX
316 changes: 316 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/complex/detail/marray.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,316 @@
//===- marray.hpp ---------------------------------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//

#pragma once

#include "common.hpp"

#include <sycl/marray.hpp>
#include <sycl/types.hpp>

namespace sycl {
inline namespace _V1 {

template <typename T> using marray_data = sycl::detail::vec_helper<T>;

template <typename T>
using marray_data_t = typename detail::vec_helper<T>::RetType;

template <typename T, std::size_t NumElements>
class marray<sycl::ext::oneapi::experimental::complex<T>, NumElements> {
private:
using ComplexDataT = sycl::ext::oneapi::experimental::complex<T>;

public:
using value_type = ComplexDataT;
using reference = ComplexDataT &;
using const_reference = const ComplexDataT &;
using iterator = ComplexDataT *;
using const_iterator = const ComplexDataT *;

private:
value_type MData[NumElements];

template <size_t... Is>
constexpr marray(const std::array<value_type, NumElements> &Arr,
std::index_sequence<Is...>)
: MData{Arr[Is]...} {}

/// FIXME: If the subscript operator is made constexpr this can be removed.
// detail::HelperFlattenMArrayArg::MArrayToArray needs to have access to
// MData.
friend class detail::HelperFlattenMArrayArg;

public:
constexpr marray() : MData{} {};

explicit constexpr marray(const value_type &arg)
: marray{sycl::detail::RepeatValue<NumElements>(
static_cast<marray_data_t<value_type>>(arg)),
std::make_index_sequence<NumElements>()} {}

template <
typename... ArgTN,
typename = std::enable_if_t<
sycl::detail::AllSuitableArgTypes<value_type, ArgTN...>::value &&
sycl::detail::GetMArrayArgsSize<ArgTN...>::value == NumElements>>
constexpr marray(const ArgTN &...Args)
: marray{
sycl::detail::MArrayArgArrayCreator<value_type, ArgTN...>::Create(
Args...),
std::make_index_sequence<NumElements>()} {}

constexpr marray(const marray<value_type, NumElements> &rhs) = default;
constexpr marray(marray<value_type, NumElements> &&rhs) = default;

// Available only when: NumElements == 1
template <std::size_t N = NumElements,
typename = std::enable_if_t<N == 1, value_type>>
operator value_type() const {
return MData[0];
}

static constexpr std::size_t size() noexcept { return NumElements; }

// subscript operator
reference operator[](std::size_t i) { return MData[i]; }
const_reference operator[](std::size_t i) const { return MData[i]; }

marray &operator=(const marray<value_type, NumElements> &rhs) = default;
marray &operator=(const value_type &rhs) {
for (std::size_t i = 0; i < NumElements; ++i) {
MData[i] = rhs;
}
return *this;
}

// iterator functions
iterator begin() { return MData; }
const_iterator begin() const { return MData; }

iterator end() { return MData + NumElements; }
const_iterator end() const { return MData + NumElements; }

#ifdef MARRAY_CPLX_OP
#error "Multiple definition of MARRAY_CPLX_OP"
#endif

// MARRAY_CPLX_OP is: +, -, *, /
#define MARRAY_CPLX_OP(op) \
friend marray operator op(const marray &lhs, const marray &rhs) { \
marray rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = lhs[i] op rhs[i]; \
} \
return rtn; \
} \
\
friend marray operator op(const marray &lhs, const value_type &rhs) { \
marray rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = lhs[i] op rhs; \
} \
return rtn; \
} \
\
friend marray operator op(const value_type &lhs, const marray &rhs) { \
marray rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = lhs op rhs[i]; \
} \
return rtn; \
}

MARRAY_CPLX_OP(+)
MARRAY_CPLX_OP(-)
MARRAY_CPLX_OP(*)
MARRAY_CPLX_OP(/)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: %
friend marray operator%(const marray &lhs, const marray &rhs) = delete;
friend marray operator%(const marray &lhs, const value_type &rhs) = delete;
friend marray operator%(const value_type &lhs, const marray &rhs) = delete;

// MARRAY_CPLX_OP is: +=, -=, *=, /=
#define MARRAY_CPLX_OP(op) \
friend marray &operator op(marray & lhs, const marray & rhs) { \
for (std::size_t i = 0; i < NumElements; ++i) { \
lhs[i] op rhs[i]; \
} \
return lhs; \
} \
\
friend marray &operator op(marray & lhs, const value_type & rhs) { \
for (std::size_t i = 0; i < NumElements; ++i) { \
lhs[i] op rhs; \
} \
return lhs; \
} \
friend marray &operator op(value_type & lhs, const marray & rhs) { \
for (std::size_t i = 0; i < NumElements; ++i) { \
lhs[i] op rhs; \
} \
return lhs; \
}

MARRAY_CPLX_OP(+=)
MARRAY_CPLX_OP(-=)
MARRAY_CPLX_OP(*=)
MARRAY_CPLX_OP(/=)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: %=
friend marray &operator%=(marray &lhs, const marray &rhs) = delete;
friend marray &operator%=(marray &lhs, const value_type &rhs) = delete;
friend marray &operator%=(value_type &lhs, const marray &rhs) = delete;

// MARRAY_CPLX_OP is: ++, --
#define MARRAY_CPLX_OP(op) \
friend marray operator op(marray &lhs, int) = delete; \
friend marray &operator op(marray & rhs) = delete;

MARRAY_CPLX_OP(++)
MARRAY_CPLX_OP(--)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: unary +, unary -
#define MARRAY_CPLX_OP(op) \
friend marray<value_type, NumElements> operator op( \
const marray<value_type, NumElements> &rhs) { \
marray<value_type, NumElements> rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = op rhs[i]; \
} \
return rtn; \
}

MARRAY_CPLX_OP(+)
MARRAY_CPLX_OP(-)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: &, |, ^
#define MARRAY_CPLX_OP(op) \
friend marray operator op(const marray &lhs, const marray &rhs) = delete; \
friend marray operator op(const marray &lhs, const value_type &rhs) = delete;

MARRAY_CPLX_OP(&)
MARRAY_CPLX_OP(|)
MARRAY_CPLX_OP(^)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: &=, |=, ^=
#define MARRAY_CPLX_OP(op) \
friend marray &operator op(marray & lhs, const marray & rhs) = delete; \
friend marray &operator op(marray & lhs, const value_type & rhs) = delete; \
friend marray &operator op(value_type & lhs, const marray & rhs) = delete;

MARRAY_CPLX_OP(&=)
MARRAY_CPLX_OP(|=)
MARRAY_CPLX_OP(^=)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: &&, ||
#define MARRAY_CPLX_OP(op) \
friend marray<bool, NumElements> operator op(const marray & lhs, \
const marray & rhs) = delete; \
friend marray<bool, NumElements> operator op( \
const marray & lhs, const value_type & rhs) = delete; \
friend marray<bool, NumElements> operator op(const value_type & lhs, \
const marray & rhs) = delete;

MARRAY_CPLX_OP(&&)
MARRAY_CPLX_OP(||)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: <<, >>
#define MARRAY_CPLX_OP(op) \
friend marray operator op(const marray &lhs, const marray &rhs) = delete; \
friend marray operator op(const marray &lhs, const value_type &rhs) = \
delete; \
friend marray operator op(const value_type &lhs, const marray &rhs) = delete;

MARRAY_CPLX_OP(<<)
MARRAY_CPLX_OP(>>)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: <<=, >>=
#define MARRAY_CPLX_OP(op) \
friend marray &operator op(marray & lhs, const marray & rhs) = delete; \
friend marray &operator op(marray & lhs, const value_type & rhs) = delete;

MARRAY_CPLX_OP(<<=)
MARRAY_CPLX_OP(>>=)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: ==, !=
#define MARRAY_CPLX_OP(op) \
friend marray<bool, NumElements> operator op(const marray & lhs, \
const marray & rhs) { \
marray<bool, NumElements> rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = lhs[i] op rhs[i]; \
} \
return rtn; \
} \
\
friend marray<bool, NumElements> operator op(const marray & lhs, \
const value_type & rhs) { \
marray<bool, NumElements> rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = lhs[i] op rhs; \
} \
return rtn; \
} \
\
friend marray<bool, NumElements> operator op(const value_type & lhs, \
const marray & rhs) { \
marray<bool, NumElements> rtn; \
for (std::size_t i = 0; i < NumElements; ++i) { \
rtn[i] = lhs op rhs[i]; \
} \
return rtn; \
}

MARRAY_CPLX_OP(==)
MARRAY_CPLX_OP(!=)

#undef MARRAY_CPLX_OP

// MARRAY_CPLX_OP is: <, >, <=, >=
#define MARRAY_CPLX_OP(op) \
friend marray<bool, NumElements> operator op(const marray & lhs, \
const marray & rhs) = delete; \
friend marray<bool, NumElements> operator op( \
const marray & lhs, const value_type & rhs) = delete; \
friend marray<bool, NumElements> operator op(const value_type & lhs, \
const marray & rhs) = delete;

MARRAY_CPLX_OP(<);
MARRAY_CPLX_OP(>);
MARRAY_CPLX_OP(<=);
MARRAY_CPLX_OP(>=);

#undef MARRAY_CPLX_OP

friend marray operator~(const marray &v) = delete;

friend marray<bool, NumElements> operator!(const marray &v) = delete;
};

} // namespace _V1
} // namespace sycl
Loading
Loading