diff --git a/src/core/grid.hh b/src/core/grid.hh index 8361b3f..74a37b9 100644 --- a/src/core/grid.hh +++ b/src/core/grid.hh @@ -1,188 +1,184 @@ /** * @file * * @author Lucas Frérot * * @section LICENSE * * Copyright (©) 2016 EPFL (Ecole Polytechnique Fédérale de * Lausanne) Laboratory (LSMS - Laboratoire de Simulation en Mécanique des * Solides) * * Tamaas is free software: you can redistribute it and/or modify it under the * terms of the GNU Lesser General Public License as published by the Free * Software Foundation, either version 3 of the License, or (at your option) any * later version. * * Tamaas is distributed in the hope that it will be useful, but WITHOUT ANY * WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR * A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more * details. * * You should have received a copy of the GNU Lesser General Public License * along with Tamaas. If not, see . * */ /* -------------------------------------------------------------------------- */ #ifndef __GRID_HH__ #define __GRID_HH__ /* -------------------------------------------------------------------------- */ #include "grid_base.hh" #include "tamaas.hh" #include #include #include #include /* -------------------------------------------------------------------------- */ __BEGIN_TAMAAS__ /** * @brief Multi-dimensional & multi-component array class * * This class is a container for multi-component data stored on a multi- * dimensional grid. * * The access function is the parenthesis operator. For a grid of dimension d, * the operator takes d+1 arguments: the first d arguments are the position on * the grid and the last one is the component of the stored data. * * It is also possible to use the access operator with only one argument, it is * then considering the grid as a flat array, accessing the given cell of the * array. */ template class Grid : public GridBase { public: /* ------------------------------------------------------------------------ */ /* Types */ /* ------------------------------------------------------------------------ */ using value_type = T; static constexpr UInt dimension = dim; /* ------------------------------------------------------------------------ */ /* Constructors */ /* ------------------------------------------------------------------------ */ public: /// Constructor by default (empty array) Grid(); /// Constructor Grid(const std::array& n, UInt nb_components); /// Wrap on given data Grid(const std::array& n, UInt nb_components, value_type * data); /// Constructor with vectors Grid(const std::vector& n, UInt nb_components); /// Constructor with initializer list Grid(const std::initializer_list& n, UInt nb_components); /// Copy constructor Grid(const Grid& o); /// Move constructor (transfers data ownership) Grid(Grid&& o); /// Destructor ~Grid() override = default; private: /// Init from standard container template void init(const Container& n, UInt nb_components); public: /* ------------------------------------------------------------------------ */ /* Common operations */ /* ------------------------------------------------------------------------ */ /// Resize array void resize(const std::array& n); /// Compute size inline UInt computeSize() const; /// Get grid dimension inline UInt getDimension() const override { return dim; } /// Compute strides virtual void computeStrides(); /// Print virtual void printself(std::ostream& str) const; /// Get sizes const std::array& sizes() const { return n; } /// Get strides const std::array& getStrides() const { return this->strides; } /* ------------------------------------------------------------------------ */ /* Access operators */ /* ------------------------------------------------------------------------ */ /// Variadic access operator (non-const) template inline T& operator()(T1... args); /// Variadic access operator template inline const T& operator()(T1... args) const; /// Tuple index access operator inline T& operator()(std::array tuple); inline const T& operator()(std::array tuple) const; private: /// Unpacking the arguments of variadic () template inline UInt unpackOffset(UInt offset, UInt index_pos, UInt index, T1... rest) const; /// End case for recursion template inline UInt unpackOffset(UInt offset, UInt index_pos, UInt index) const; /// Computing offset for a tuple index inline UInt computeOffset(std::array tuple) const; /* ------------------------------------------------------------------------ */ /* Move/Copy operators */ /* ------------------------------------------------------------------------ */ public: + using GridBase::operator=; // = operator Grid& operator=(const Grid& other); // = operator (not const input, otherwise gcc is confused) Grid& operator=(Grid& other); // = operator (move) Grid& operator=(Grid&& other); - // = operator for scalar - Grid& operator=(T val) { - GridBase::operator=(val); - return *this; - } // Copy data from another grid template void copy(const Grid& other); // Move data from another grid template void move(Grid&& other); // Wrap memory (non-const) void wrap(Grid & other) { GridBase::wrap(other); std::copy(other.n.begin(), other.n.end(), n.begin()); this->computeStrides(); } // Wrap memory void wrap(const Grid & other) { GridBase::wrap(other); std::copy(other.n.begin(), other.n.end(), n.begin()); this->computeStrides(); } /* ------------------------------------------------------------------------ */ /* Member variables */ /* ------------------------------------------------------------------------ */ protected: std::array n; ///< shape of grid: size per dimension std::array strides; ///< strides for access }; __END_TAMAAS__ /* -------------------------------------------------------------------------- */ /* Inline/template function definitions */ /* -------------------------------------------------------------------------- */ #include "grid_tmpl.hh" /* -------------------------------------------------------------------------- */ #endif // __GRID_HH__ diff --git a/src/core/grid_base.hh b/src/core/grid_base.hh index 0d17874..bcd296b 100644 --- a/src/core/grid_base.hh +++ b/src/core/grid_base.hh @@ -1,298 +1,338 @@ /** * @file * * @author Lucas Frérot * * @section LICENSE * * Copyright (©) 2016 EPFL (Ecole Polytechnique Fédérale de * Lausanne) Laboratory (LSMS - Laboratoire de Simulation en Mécanique des * Solides) * * Tamaas is free software: you can redistribute it and/or modify it under the * terms of the GNU Lesser General Public License as published by the Free * Software Foundation, either version 3 of the License, or (at your option) any * later version. * * Tamaas is distributed in the hope that it will be useful, but WITHOUT ANY * WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR * A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more * details. * * You should have received a copy of the GNU Lesser General Public License * along with Tamaas. If not, see . * */ /* -------------------------------------------------------------------------- */ #ifndef __GRID_BASE_HH__ #define __GRID_BASE_HH__ /* -------------------------------------------------------------------------- */ #include "array.hh" #include "iterator.hh" #include "loop.hh" #include "tamaas.hh" +#include "static_types.hh" #include #include #include #include /* -------------------------------------------------------------------------- */ __BEGIN_TAMAAS__ /** * @brief Dimension agnostic grid with components stored per points */ template class GridBase { public: /// Constructor by default GridBase() = default; /// Copy constructor GridBase(const GridBase& o) : data(o.data), nb_components(o.nb_components), offset(o.offset) {} /// Move constructor (transfers data ownership) GridBase(GridBase&& o) : data(std::move(o.data)), nb_components(std::exchange(o.nb_components, 1)), offset(std::exchange(o.offset, 0)) {} /// Destructor virtual ~GridBase() = default; /* ------------------------------------------------------------------------ */ /* Iterator class */ /* ------------------------------------------------------------------------ */ public: using iterator = iterator_::iterator; using const_iterator = iterator_::iterator; public: /// Get grid dimension virtual UInt getDimension() const { return 1; } /// Get internal data pointer (const) const T* getInternalData() const { return this->data.data(); } /// Get internal data pointer (non-const) T* getInternalData() { return this->data.data(); } /// Get number of components UInt getNbComponents() const { return nb_components; } /// Set number of components void setNbComponents(UInt n) { nb_components = n; } /// Get offset UInt getOffset() const { return offset; } /// Get total size virtual UInt dataSize() const { return this->data.size(); } /// Get number of points UInt getNbPoints() const { return this->dataSize() / this->getNbComponents(); } /// Set components void uniformSetComponents(const GridBase& vec); /* ------------------------------------------------------------------------ */ /* Iterator methods */ /* ------------------------------------------------------------------------ */ virtual iterator begin(UInt n = 1) { return iterator(this->getInternalData() + this->offset, 0, n); } virtual iterator end(UInt n = 1) { return iterator(this->getInternalData() + this->offset, this->dataSize(), n); } virtual const_iterator begin(UInt n = 1) const { return const_iterator(this->getInternalData() + this->offset, 0, n); } virtual const_iterator end(UInt n = 1) const { return const_iterator(this->getInternalData() + this->offset, this->dataSize(), n); } /* ------------------------------------------------------------------------ */ /* Operators */ /* ------------------------------------------------------------------------ */ inline T& operator()(UInt i) { return this->data[i]; } inline const T& operator()(UInt i) const { return this->data[i]; } #define VEC_OPERATOR(op) \ template \ GridBase& operator op(const GridBase& other) VEC_OPERATOR(+=); VEC_OPERATOR(*=); VEC_OPERATOR(-=); VEC_OPERATOR(/=); #undef VEC_OPERATOR template T dot(const GridBase& other) const; #define SCALAR_OPERATOR(op) GridBase& operator op(const T& e) SCALAR_OPERATOR(+=); SCALAR_OPERATOR(*=); SCALAR_OPERATOR(-=); SCALAR_OPERATOR(/=); SCALAR_OPERATOR(=); #undef SCALAR_OPERATOR - GridBase& operator=(const GridBase& o); +/// Broadcast operators +#define BROADCAST_OPERATOR(op) \ + template \ + GridBase& operator op(const StaticArray& b) + + BROADCAST_OPERATOR(+=); + BROADCAST_OPERATOR(-=); + BROADCAST_OPERATOR(*=); + BROADCAST_OPERATOR(/=); + BROADCAST_OPERATOR(=); +#undef BROADCAST_OPERATOR /* ------------------------------------------------------------------------ */ /* Min/max */ /* ------------------------------------------------------------------------ */ T min() const; T max() const; T sum() const; T mean() const { return sum() / static_cast(dataSize()); } T var() const; /* ------------------------------------------------------------------------ */ /* Move/Copy */ /* ------------------------------------------------------------------------ */ public: + GridBase& operator=(const GridBase& o) { + this->copy(o); + return *this; + } + + GridBase& operator=(GridBase& o) { + this->copy(o); + return *this; + } + + GridBase& operator=(GridBase&& o) { + this->move(std::move(o)); + return *this; + } + template void copy(const GridBase& other) { data = other.data; nb_components = other.nb_components; offset = other.offset; } template void move(GridBase&& other) { data = std::move(other.data); nb_components = std::exchange(other.nb_components, 1); offset = std::exchange(other.offset, 0); } GridBase& wrap(GridBase& o) { data.wrap(o.data); nb_components = o.nb_components; offset = o.offset; return *this; } GridBase& wrap(const GridBase& o) { data.wrap(o.data); nb_components = o.nb_components; offset = o.offset; return *this; } protected: Array data; UInt nb_components = 1; UInt offset = 0; }; /* -------------------------------------------------------------------------- */ template void GridBase::uniformSetComponents(const GridBase& vec) { if (!(vec.dataSize() == this->nb_components)) TAMAAS_EXCEPTION("Cannot set grid field with values of vector"); auto begin_it(begin()); auto end_it(end()); // silencing nvcc warnings #pragma omp parallel for for (auto it = begin(); it < end_it; ++it) { UInt i = it - begin_it; *it = vec.data[i % this->nb_components]; } } /* -------------------------------------------------------------------------- */ #define SCALAR_OPERATOR_IMPL(op) \ template \ inline GridBase& GridBase::operator op(const T& e) { \ Loop::loop([e] CUDA_LAMBDA(T& val) { val op e; }, *this); \ return *this; \ } SCALAR_OPERATOR_IMPL(+=); SCALAR_OPERATOR_IMPL(*=); SCALAR_OPERATOR_IMPL(-=); SCALAR_OPERATOR_IMPL(/=); SCALAR_OPERATOR_IMPL(=); #undef SCALAR_OPERATOR_IMPL /* -------------------------------------------------------------------------- */ #define VEC_OPERATOR_IMPL(op) \ template \ template \ inline GridBase& GridBase::operator op(const GridBase& other) { \ TAMAAS_ASSERT(other.dataSize() == this->dataSize(), \ "surface size does not match"); \ Loop::loop([] CUDA_LAMBDA(T& x, const T1& y) { x op y; }, *this, other); \ return *this; \ } VEC_OPERATOR_IMPL(+=); VEC_OPERATOR_IMPL(-=); VEC_OPERATOR_IMPL(*=); VEC_OPERATOR_IMPL(/=); #undef VEC_OPERATOR /* -------------------------------------------------------------------------- */ -template -GridBase& GridBase::operator=(const GridBase& o) { - this->copy(o); - return *this; -} +#define BROADCAST_OPERATOR_IMPL(op) \ + template \ + template \ + GridBase& GridBase::operator op(const StaticArray& b) { \ + TAMAAS_ASSERT(this->dataSize() % size == 0, \ + "Broadcast operator cannot broadcast" \ + << size << " on array of size " << this->dataSize()); \ + Tensor a(b); \ + Loop::stridedLoop([a] CUDA_LAMBDA(VectorProxy&& g) { g op a; }, \ + *this); \ + return *this; \ + } -/* -------------------------------------------------------------------------- */ +BROADCAST_OPERATOR_IMPL(+=); +BROADCAST_OPERATOR_IMPL(-=); +BROADCAST_OPERATOR_IMPL(*=); +BROADCAST_OPERATOR_IMPL(/=); +BROADCAST_OPERATOR_IMPL(=); +#undef BROADCAST_OPERATOR_IMPL -using op = operation; +/* -------------------------------------------------------------------------- */ template inline T GridBase::min() const { // const auto id = [] CUDA_LAMBDA (const T&x){return x;}; - return Loop::reduce([] CUDA_LAMBDA(const T& x) { return x; }, *this); + return Loop::reduce([] CUDA_LAMBDA(const T& x) { return x; }, + *this); } /* -------------------------------------------------------------------------- */ template inline T GridBase::max() const { - return Loop::reduce([] CUDA_LAMBDA(const T& x) { return x; }, *this); + return Loop::reduce([] CUDA_LAMBDA(const T& x) { return x; }, + *this); } /* -------------------------------------------------------------------------- */ template inline T GridBase::sum() const { - return Loop::reduce([] CUDA_LAMBDA(const T& x) { return x; }, - *this); + return Loop::reduce([] CUDA_LAMBDA(const T& x) { return x; }, + *this); } /* -------------------------------------------------------------------------- */ template inline T GridBase::var() const { const T mu = mean(); - const T var = Loop::reduce( + const T var = Loop::reduce( [mu] CUDA_LAMBDA(const T& x) { return (x - mu) * (x - mu); }, *this); return var / (dataSize() - 1); } /* -------------------------------------------------------------------------- */ template template T GridBase::dot(const GridBase& other) const { - return Loop::reduce( + return Loop::reduce( [] CUDA_LAMBDA(const T& x, const T1& y) { return x * y; }, *this, other); } __END_TAMAAS__ #endif // __GRID_BASE_HH__ diff --git a/src/core/grid_tmpl.hh b/src/core/grid_tmpl.hh index 2f17992..922d217 100644 --- a/src/core/grid_tmpl.hh +++ b/src/core/grid_tmpl.hh @@ -1,170 +1,170 @@ /** * @file * * @author Lucas Frérot * * @section LICENSE * * Copyright (©) 2016 EPFL (Ecole Polytechnique Fédérale de * Lausanne) Laboratory (LSMS - Laboratoire de Simulation en Mécanique des * Solides) * * Tamaas is free software: you can redistribute it and/or modify it under the * terms of the GNU Lesser General Public License as published by the Free * Software Foundation, either version 3 of the License, or (at your option) any * later version. * * Tamaas is distributed in the hope that it will be useful, but WITHOUT ANY * WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR * A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more * details. * * You should have received a copy of the GNU Lesser General Public License * along with Tamaas. If not, see . * */ /* -------------------------------------------------------------------------- */ #ifndef __GRID_TMPL_HH__ #define __GRID_TMPL_HH__ /* -------------------------------------------------------------------------- */ #include "grid.hh" #include "tamaas.hh" __BEGIN_TAMAAS__ template template void Grid::init(const T1& n, UInt nb_components) { std::copy(n.begin(), n.end(), this->n.begin()); this->nb_components = nb_components; this->resize(this->n); } /* -------------------------------------------------------------------------- */ template inline UInt Grid::computeSize() const { UInt size = 1; for (UInt i = 0; i < dim; i++) size *= n[i]; size *= this->nb_components; return size; } /* -------------------------------------------------------------------------- */ template template inline UInt Grid::unpackOffset(UInt offset, UInt index_pos, UInt index, T1... rest) const { offset += index * strides[index_pos]; return unpackOffset(offset, index_pos + 1, rest...); // tail-rec bb } template template inline UInt Grid::unpackOffset(UInt offset, UInt index_pos, UInt index) const { return offset + index * strides[index_pos]; } template template inline T& Grid::operator()(T1... args) { /// Checking access integrity constexpr UInt nargs = sizeof...(T1); static_assert(nargs == dim + 1 || nargs == 1 || nargs == dim, "number of arguments in operator() does not match dimension"); constexpr UInt start = (nargs == 1) ? dim : 0; UInt offset = unpackOffset(this->offset, start, args...); return this->data[offset]; } template template inline const T& Grid::operator()(T1... args) const { /// Checking access integrity constexpr UInt nargs = sizeof...(T1); static_assert(nargs == dim + 1 || nargs == 1 || nargs == dim, "number of arguments in operator() does not match dimension"); constexpr UInt start = (nargs == 1) ? dim : 0; UInt offset = unpackOffset(this->offset, start, args...); return this->data[offset]; } /* -------------------------------------------------------------------------- */ template inline UInt Grid::computeOffset(std::array tuple) const { UInt offset = this->offset; return std::inner_product(tuple.begin(), tuple.end(), strides.begin(), offset); } template inline T& Grid::operator()(std::array tuple) { UInt offset = computeOffset(tuple); return this->data[offset]; } template inline const T& Grid:: operator()(std::array tuple) const { UInt offset = computeOffset(tuple); return this->data[offset]; } /* -------------------------------------------------------------------------- */ template Grid& Grid::operator=(const Grid& other) { this->copy(other); return *this; } template Grid& Grid::operator=(Grid& other) { this->copy(other); return *this; } template Grid& Grid::operator=(Grid&& other) { - this->move(std::forward>(other)); + this->move(std::move(other)); return *this; } /* -------------------------------------------------------------------------- */ template template void Grid::copy(const Grid& other) { GridBase::copy(other); this->n = other.n; this->strides = other.strides; } template template void Grid::move(Grid&& other) { GridBase::move(std::forward>(other)); this->n = std::move(other.n); this->strides = std::move(other.strides); } /* -------------------------------------------------------------------------- */ /* Stream operator */ /* -------------------------------------------------------------------------- */ template inline std::ostream& operator<<(std::ostream& stream, const Grid& _this) { _this.printself(stream); return stream; } __END_TAMAAS__ /* -------------------------------------------------------------------------- */ #endif // __GRID_TMPL_HH__ diff --git a/src/core/static_types.hh b/src/core/static_types.hh index 303093f..f15e0f3 100644 --- a/src/core/static_types.hh +++ b/src/core/static_types.hh @@ -1,328 +1,333 @@ /** * @file * * @author Lucas Frérot * * @section LICENSE * * Copyright (©) 2017 EPFL (Ecole Polytechnique Fédérale de * Lausanne) Laboratory (LSMS - Laboratoire de Simulation en Mécanique des * Solides) * * Tamaas is free software: you can redistribute it and/or modify it under the * terms of the GNU Lesser General Public License as published by the Free * Software Foundation, either version 3 of the License, or (at your option) any * later version. * * Tamaas is distributed in the hope that it will be useful, but WITHOUT ANY * WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR * A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more * details. * * You should have received a copy of the GNU Lesser General Public License * along with Tamaas. If not, see . * */ /* -------------------------------------------------------------------------- */ #ifndef __STATIC_TYPES_HH__ #define __STATIC_TYPES_HH__ /* -------------------------------------------------------------------------- */ #include "tamaas.hh" __BEGIN_TAMAAS__ /* -------------------------------------------------------------------------- */ namespace detail { template struct product_tail_rec : product_tail_rec {}; template struct product_tail_rec : std::integral_constant {}; template struct get_rec : get_rec {}; template struct get_rec<0, n, ns...> : std::integral_constant {}; } template struct product : detail::product_tail_rec<1, ns...> {}; template struct get : detail::get_rec {}; template struct is_arithmetic : std::is_arithmetic {}; template struct is_arithmetic> : std::true_type {}; /** * @brief Static Array * * This class is meant to be a small and fast object for intermediate * calculations, possibly on wrapped memory belonging to a grid. Support type * show be either a pointer or a C array. It should not contain any virtual * method. */ template class StaticArray { static_assert(std::is_array::value || std::is_pointer::value, "the support type of StaticArray should be either a pointer or " "a C-array"); using T = DataType; public: /// Access operator __device__ __host__ T& operator()(UInt i) { // TAMAAS_ASSERT(i < n, "Access out of bounds"); return _mem[i]; } /// Access operator __device__ __host__ const T& operator()(UInt i) const { // TAMAAS_ASSERT(i < n, "Access out of bounds"); return _mem[i]; } /// L2 norm squared __device__ __host__ T l2squared() const { T res = 0; for (UInt i = 0; i < size; ++i) - res += _mem[i] * _mem[i]; + res += (*this)(i) * (*this)(i); return res; } /// L2 norm __device__ __host__ T l2norm() const { return std::sqrt(l2squared()); } #define VECTOR_OP(op) \ template \ __device__ __host__ void operator op(const StaticArray& o) { \ for (UInt i = 0; i < size; ++i) \ (*this)(i) op o(i); \ } VECTOR_OP(+=) VECTOR_OP(-=) VECTOR_OP(*=) VECTOR_OP(/=) #undef VECTOR_OP #define SCALAR_OP(op) \ template \ __device__ __host__ \ typename std::enable_if::value, StaticArray&>::type \ operator op(const T1& x) { \ for (UInt i = 0; i < size; ++i) \ (*this)(i) op x; \ return *this; \ } SCALAR_OP(+=) SCALAR_OP(-=) SCALAR_OP(*=) SCALAR_OP(/=) SCALAR_OP(=) #undef SCALAR_OP - /// Overriding the implit copy operator + /// Overriding the implicit copy operator StaticArray& operator=(const StaticArray& o) { return this->copy(o); } + template + void operator=(const StaticArray& o) { + this->copy(o); + } + template StaticArray& copy(const StaticArray& o) { for (UInt i = 0; i < size; ++i) (*this)(i) = o(i); return *this; } protected: SupportType _mem; }; /** * @brief Static Tensor * * This class implements a multi-dimensional tensor behavior. */ template class StaticTensor : public StaticArray::value> { using parent = StaticArray::value>; using T = DataType; public: static constexpr UInt dim = sizeof...(dims); using parent::operator=; private: template static UInt unpackOffset(UInt offset, UInt index, Idx... rest) { constexpr UInt size = sizeof...(rest); offset += index; offset *= get::value; return unpackOffset(offset, rest...); } template static UInt unpackOffset(UInt offset, UInt index) { return offset + index; } public: template const T& operator()(Idx... idx) const { - return this->_mem[unpackOffset(0, idx...)]; + return parent::operator()(unpackOffset(0, idx...)); } template T& operator()(Idx... idx) { - return this->_mem[unpackOffset(0, idx...)]; + return parent::operator()(unpackOffset(0, idx...)); } }; /* -------------------------------------------------------------------------- */ /* Common Static Types */ /* -------------------------------------------------------------------------- */ template using StaticMatrix = StaticTensor; /// Vector class with size determined at compile-time template class StaticVector : public StaticTensor { using T = DataType; public: using StaticTensor::operator=; /// Matrix-vector product template __device__ __host__ void mul(const StaticMatrix& mat, const StaticVector& vec) { *this = T(0); for (UInt i = 0; i < n; ++i) for (UInt j = 0; j < m; ++j) (*this)(i) += mat(i, j) * vec(j); } }; /* -------------------------------------------------------------------------- */ /* Proxy Static Types */ /* -------------------------------------------------------------------------- */ /// Proxy type for tensor template