From 5eebd1e4bfce1a921409443a3662a1aeabd4ac6e Mon Sep 17 00:00:00 2001 From: Ilya Burylov Date: Wed, 23 Dec 2020 00:39:54 +0300 Subject: [PATCH] [SYCL] Add marray class as defined by SYCL 2020 provisional (#2928) Signed-off-by: iburylov --- sycl/include/CL/sycl.hpp | 1 + sycl/include/CL/sycl/marray.hpp | 349 ++++++++++++++++++++++++ sycl/test/basic_tests/marray/marray.cpp | 87 ++++++ 3 files changed, 437 insertions(+) create mode 100755 sycl/include/CL/sycl/marray.hpp create mode 100755 sycl/test/basic_tests/marray/marray.cpp diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp index e90d631a13585..e1658dc33493b 100644 --- a/sycl/include/CL/sycl.hpp +++ b/sycl/include/CL/sycl.hpp @@ -33,6 +33,7 @@ #include #include #include +#include #include #include #include diff --git a/sycl/include/CL/sycl/marray.hpp b/sycl/include/CL/sycl/marray.hpp new file mode 100755 index 0000000000000..7e2a4baf930cc --- /dev/null +++ b/sycl/include/CL/sycl/marray.hpp @@ -0,0 +1,349 @@ +//==----------------- marray.hpp --- Implements marray classes -------------==// +// +// 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 +#include +#include +#include +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { + +/// Provides a cross-patform math array class template that works on +/// SYCL devices as well as in host C++ code. +/// +/// \ingroup sycl_api +template class marray { + using DataT = Type; + +public: + using value_type = Type; + using reference = Type &; + using const_reference = const Type &; + using iterator = Type *; + using const_iterator = const Type *; + +private: + value_type MData[NumElements]; + + template struct conjunction : std::true_type {}; + template + struct conjunction + : std::conditional, B1>::type {}; + + // TypeChecker is needed for (const ArgTN &... Args) ctor to validate Args. + template + struct TypeChecker : std::is_convertible {}; + + // Shortcuts for Args validation in (const ArgTN &... Args) ctor. + template + using EnableIfSuitableTypes = typename std::enable_if< + conjunction...>::value>::type; + +public: + marray() : MData{} {} + + explicit marray(const Type &Arg) { + for (std::size_t I = 0; I < NumElements; ++I) { + MData[I] = Arg; + } + } + + template < + typename... ArgTN, typename = EnableIfSuitableTypes, + typename = typename std::enable_if::type> + marray(const ArgTN &... Args) : MData{Args...} {} + + marray(const marray &Rhs) { + for (std::size_t I = 0; I < NumElements; ++I) { + MData[I] = Rhs.MData[I]; + } + } + + marray(marray &&Rhs) { + for (std::size_t I = 0; I < NumElements; ++I) { + MData[I] = Rhs.MData[I]; + } + } + + // Available only when: NumElements == 1 + template > + operator Type() const { + return MData[0]; + } + + static constexpr std::size_t size() noexcept { return NumElements; } + + // subscript operator + reference operator[](std::size_t index) { return MData[index]; } + + const_reference operator[](std::size_t index) const { return MData[index]; } + + marray &operator=(const marray &Rhs) { + for (std::size_t I = 0; I < NumElements; ++I) { + MData[I] = Rhs.MData[I]; + } + return *this; + } + + // broadcasting operator + marray &operator=(const 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 __SYCL_BINOP +#error "Undefine __SYCL_BINOP macro" +#endif + +#ifdef __SYCL_BINOP_INTEGRAL +#error "Undefine __SYCL_BINOP_INTEGRAL macro" +#endif + +#define __SYCL_BINOP(BINOP, OPASSIGN) \ + friend marray operator BINOP(const marray &Lhs, const marray &Rhs) { \ + marray Ret; \ + for (size_t I = 0; I < NumElements; ++I) { \ + Ret[I] = Lhs[I] BINOP Rhs[I]; \ + } \ + return Ret; \ + } \ + template \ + friend typename std::enable_if< \ + std::is_convertible::value && \ + (std::is_fundamental::value || \ + std::is_same::type, half>::value), \ + marray>::type \ + operator BINOP(const marray &Lhs, const T &Rhs) { \ + return Lhs BINOP marray(static_cast(Rhs)); \ + } \ + friend marray &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \ + Lhs = Lhs BINOP Rhs; \ + return Lhs; \ + } \ + template \ + friend typename std::enable_if::type operator OPASSIGN( \ + marray &Lhs, const DataT &Rhs) { \ + Lhs = Lhs BINOP marray(Rhs); \ + return Lhs; \ + } + +#define __SYCL_BINOP_INTEGRAL(BINOP, OPASSIGN) \ + template \ + friend typename std::enable_if::value, marray> \ + operator BINOP(const marray &Lhs, const marray &Rhs) { \ + marray Ret; \ + for (size_t I = 0; I < NumElements; ++I) { \ + Ret[I] = Lhs[I] BINOP Rhs[I]; \ + } \ + return Ret; \ + } \ + template \ + friend typename std::enable_if::value && \ + std::is_integral::value && \ + std::is_integral::value, \ + marray>::type \ + operator BINOP(const marray &Lhs, const T &Rhs) { \ + return Lhs BINOP marray(static_cast(Rhs)); \ + } \ + template \ + friend typename std::enable_if::value, marray> \ + &operator OPASSIGN(marray &Lhs, const marray &Rhs) { \ + Lhs = Lhs BINOP Rhs; \ + return Lhs; \ + } \ + template \ + friend typename std::enable_if::value, \ + marray &>::type \ + operator OPASSIGN(marray &Lhs, const DataT &Rhs) { \ + Lhs = Lhs BINOP marray(Rhs); \ + return Lhs; \ + } + + __SYCL_BINOP(+, +=) + __SYCL_BINOP(-, -=) + __SYCL_BINOP(*, *=) + __SYCL_BINOP(/, /=) + + __SYCL_BINOP_INTEGRAL(%, %=) + __SYCL_BINOP_INTEGRAL(|, |=) + __SYCL_BINOP_INTEGRAL(&, &=) + __SYCL_BINOP_INTEGRAL(^, ^=) + __SYCL_BINOP_INTEGRAL(>>, >>=) + __SYCL_BINOP_INTEGRAL(<<, <<=) +#undef __SYCL_BINOP +#undef __SYCL_BINOP_INTEGRAL + +#ifdef __SYCL_RELLOGOP +#error "Undefine __SYCL_RELLOGOP macro" +#endif + +#ifdef __SYCL_RELLOGOP_INTEGRAL +#error "Undefine __SYCL_RELLOGOP_INTEGRAL macro" +#endif + +#define __SYCL_RELLOGOP(RELLOGOP) \ + friend marray operator RELLOGOP(const marray &Lhs, \ + const marray &Rhs) { \ + marray Ret; \ + for (size_t I = 0; I < NumElements; ++I) { \ + Ret[I] = Lhs[I] RELLOGOP Rhs[I]; \ + } \ + return Ret; \ + } \ + template \ + friend typename std::enable_if::value && \ + (std::is_fundamental::value || \ + std::is_same::value), \ + marray>::type \ + operator RELLOGOP(const marray &Lhs, const T &Rhs) { \ + return Lhs RELLOGOP marray(static_cast(Rhs)); \ + } + +#define __SYCL_RELLOGOP_INTEGRAL(RELLOGOP) \ + template \ + friend typename std::enable_if::value, \ + marray>::type \ + operator RELLOGOP(const marray &Lhs, const marray &Rhs) { \ + marray Ret; \ + for (size_t I = 0; I < NumElements; ++I) { \ + Ret[I] = Lhs[I] RELLOGOP Rhs[I]; \ + } \ + return Ret; \ + } \ + template \ + friend typename std::enable_if::value && \ + std::is_integral::value && \ + std::is_integral::value, \ + marray>::type \ + operator RELLOGOP(const marray &Lhs, const T &Rhs) { \ + return Lhs RELLOGOP marray(static_cast(Rhs)); \ + } + + __SYCL_RELLOGOP(==) + __SYCL_RELLOGOP(!=) + __SYCL_RELLOGOP(>) + __SYCL_RELLOGOP(<) + __SYCL_RELLOGOP(>=) + __SYCL_RELLOGOP(<=) + + __SYCL_RELLOGOP_INTEGRAL(&&) + __SYCL_RELLOGOP_INTEGRAL(||) + +#undef __SYCL_RELLOGOP +#undef __SYCL_RELLOGOP_INTEGRAL + +#ifdef __SYCL_UOP +#error "Undefine __SYCL_UOP macro" +#endif + +#define __SYCL_UOP(UOP, OPASSIGN) \ + friend marray &operator UOP(marray &Lhs) { \ + Lhs OPASSIGN 1; \ + return Lhs; \ + } \ + friend marray operator UOP(marray &Lhs, int) { \ + marray Ret(Lhs); \ + Lhs OPASSIGN 1; \ + return Ret; \ + } + + __SYCL_UOP(++, +=) + __SYCL_UOP(--, -=) +#undef __SYCL_UOP + + // Available only when: dataT != cl_float && dataT != cl_double + // && dataT != cl_half + template + friend typename std::enable_if::value, marray>::type + operator~(marray &Lhs) { + marray Ret; + for (size_t I = 0; I < NumElements; ++I) { + Ret[I] = ~Lhs[I]; + } + return Ret; + } + + friend marray operator!(marray &Lhs) { + marray Ret; + for (size_t I = 0; I < NumElements; ++I) { + Ret[I] = !Lhs[I]; + } + return Ret; + } + + friend marray operator+(marray &Lhs) { + marray Ret; + for (size_t I = 0; I < NumElements; ++I) { + Ret[I] = +Lhs[I]; + } + return Ret; + } + + friend marray operator-(marray &Lhs) { + marray Ret; + for (size_t I = 0; I < NumElements; ++I) { + Ret[I] = -Lhs[I]; + } + return Ret; + } +}; + +#define __SYCL_MAKE_MARRAY_ALIAS(ALIAS, TYPE, N) \ + using ALIAS##N = cl::sycl::marray; + +#define __SYCL_MAKE_MARRAY_ALIASES_FOR_ARITHMETIC_TYPES(N) \ + __SYCL_MAKE_MARRAY_ALIAS(mchar, char, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mshort, short, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mint, int, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mlong, long, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mfloat, float, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mdouble, double, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mhalf, half, N) + +#define __SYCL_MAKE_MARRAY_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES(N) \ + __SYCL_MAKE_MARRAY_ALIAS(mschar, signed char, N) \ + __SYCL_MAKE_MARRAY_ALIAS(muchar, unsigned char, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mushort, unsigned short, N) \ + __SYCL_MAKE_MARRAY_ALIAS(muint, unsigned int, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mulong, unsigned long, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mlonglong, long long, N) \ + __SYCL_MAKE_MARRAY_ALIAS(mulonglong, unsigned long long, N) + +#define __SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(N) \ + __SYCL_MAKE_MARRAY_ALIASES_FOR_ARITHMETIC_TYPES(N) \ + __SYCL_MAKE_MARRAY_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES(N) + +__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(2) +__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(3) +__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(4) +__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(8) +__SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH(16) + +#undef __SYCL_MAKE_MARRAY_ALIAS +#undef __SYCL_MAKE_MARRAY_ALIASES_FOR_ARITHMETIC_TYPES +#undef __SYCL_MAKE_MARRAY_ALIASES_FOR_SIGNED_AND_UNSIGNED_TYPES +#undef __SYCL_MAKE_MARRAY_ALIASES_FOR_MARRAY_LENGTH + +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/basic_tests/marray/marray.cpp b/sycl/test/basic_tests/marray/marray.cpp new file mode 100755 index 0000000000000..5b459d0428425 --- /dev/null +++ b/sycl/test/basic_tests/marray/marray.cpp @@ -0,0 +1,87 @@ +// RUN: %clangxx %s -o %t.out -lsycl -I %sycl_include +// RUN: %t.out +//==--------------- marray.cpp - SYCL marray test --------------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include +using namespace sycl; + +int main() { + // Constructing vector from a scalar + sycl::marray marray_from_one_elem(1); + + // Check broadcasting operator= + sycl::marray b_marray(1.0); + b_marray = 0.5; + assert(static_cast(b_marray[0]) == static_cast(0.5)); + assert(static_cast(b_marray[1]) == static_cast(0.5)); + assert(static_cast(b_marray[2]) == static_cast(0.5)); + assert(static_cast(b_marray[3]) == static_cast(0.5)); + + // Check that [u]long[n] type aliases match marray<[unsigned] long, n> types. + assert((std::is_same, sycl::mlong2>::value)); + assert((std::is_same, sycl::mlong3>::value)); + assert((std::is_same, sycl::mlong4>::value)); + assert((std::is_same, sycl::mlong8>::value)); + assert((std::is_same, sycl::mlong16>::value)); + assert((std::is_same, sycl::mulong2>::value)); + assert((std::is_same, sycl::mulong3>::value)); + assert((std::is_same, sycl::mulong4>::value)); + assert((std::is_same, sycl::mulong8>::value)); + assert( + (std::is_same, sycl::mulong16>::value)); + + mint3 t000; + mint3 t222{2}; + mint3 t123{1, 2, 3}; + mint3 tcpy{t123}; + mint3 t___; + sycl::marray b___; + + // test default ctor + assert(t000[0] == 0 && t000[1] == 0 && t000[2] == 0); + + // test constant ctor + assert(t222[0] == 2 && t222[1] == 2 && t222[2] == 2); + + // test vararg ctor + assert(t123[0] == 1 && t123[1] == 2 && t123[2] == 3); + + // test copy ctor + assert(tcpy[0] == 1 && tcpy[1] == 2 && tcpy[2] == 3); + + // test iterators + for (auto &a : t___) { + a = 9; + } + assert(t___[0] == 9 && t___[1] == 9 && t___[2] == 9); + + // test relation operator forms + t___ = t123 + t222; + assert(t___[0] == 3 && t___[1] == 4 && t___[2] == 5); + t___ = t123 - 1; + assert(t___[0] == 0 && t___[1] == 1 && t___[2] == 2); + t___ += t123; + assert(t___[0] == 1 && t___[1] == 3 && t___[2] == 5); + t___ -= 1; + assert(t___[0] == 0 && t___[1] == 2 && t___[2] == 4); + + // test unary operator forms + t___++; + assert(t___[0] == 1 && t___[1] == 3 && t___[2] == 5); + --t___; + assert(t___[0] == 0 && t___[1] == 2 && t___[2] == 4); + + // test relation operator forms + b___ = t123 > t222; + assert(b___[0] == false && b___[1] == false && b___[2] == true); + b___ = t123 < 2; + assert(b___[0] == true && b___[1] == false && b___[2] == false); + + return 0; +}