diff --git a/examples/boundaries/boundaries.cpp b/examples/boundaries/boundaries.cpp index 89185a4344..9d29e54cc9 100644 --- a/examples/boundaries/boundaries.cpp +++ b/examples/boundaries/boundaries.cpp @@ -64,11 +64,9 @@ template struct direction_bc_input { T value; - GT_FUNCTION - direction_bc_input() : value(1) {} + constexpr direction_bc_input() : value(1) {} - GT_FUNCTION - direction_bc_input(T v) : value(v) {} + constexpr direction_bc_input(T v) : value(v) {} // This is the primary template and it will be picked when all the other specializations below fail to be selected. // It is important for debugging to note that, if a needed specialization i missing, this version will be selected. diff --git a/examples/stencil/tridiagonal_solver.cpp b/examples/stencil/tridiagonal_solver.cpp index 5556c4bfc9..a6d36712db 100644 --- a/examples/stencil/tridiagonal_solver.cpp +++ b/examples/stencil/tridiagonal_solver.cpp @@ -114,8 +114,7 @@ int main() { // Here we make the computation, specifying the backend, the grid (iteration space), binding of the spec arguments // to the fields - st::run( - spec, stencil_backend_t(), grid, st::make_global_parameter(-1), st::make_global_parameter(3), sup, rhs, out); + st::run(spec, stencil_backend_t(), grid, st::global_parameter(-1), st::global_parameter(3), sup, rhs, out); // In this simple example the solution is known and we can easily check it. auto view = out->const_host_view(); diff --git a/examples/stencil_type_erasure/interpolate_stencil.cpp b/examples/stencil_type_erasure/interpolate_stencil.cpp index b984a32353..89107a4c05 100644 --- a/examples/stencil_type_erasure/interpolate_stencil.cpp +++ b/examples/stencil_type_erasure/interpolate_stencil.cpp @@ -48,6 +48,6 @@ namespace { std::function make_interpolate_stencil(grid_t grid, double weight) { return [grid = std::move(grid), weight](inputs in, outputs out) { run_single_stage( - interpolate_stage(), stencil_backend_t(), grid, in.in1, in.in2, make_global_parameter(weight), out.out); + interpolate_stage(), stencil_backend_t(), grid, in.in1, in.in2, global_parameter(weight), out.out); }; } diff --git a/include/gridtools/boundaries/apply_gpu.hpp b/include/gridtools/boundaries/apply_gpu.hpp index 66cb731f02..cd3541492a 100644 --- a/include/gridtools/boundaries/apply_gpu.hpp +++ b/include/gridtools/boundaries/apply_gpu.hpp @@ -69,7 +69,7 @@ namespace gridtools { private: /** Computation of the bit-index in `m_predicate_values` of the given direction. */ template - GT_FUNCTION static constexpr uint_t direction_index(direction) { + static constexpr uint_t direction_index(direction) { // computation of the bit-index of the given direction constexpr int_t stride_i = 9; constexpr int_t stride_j = 3; @@ -133,23 +133,12 @@ namespace gridtools { } } - GT_FUNCTION - uint_t max() const { return m_sorted[0]; } - - GT_FUNCTION - uint_t min() const { return m_sorted[2]; } - - GT_FUNCTION - uint_t median() const { return m_sorted[1]; } - - GT_FUNCTION - uint_t size(uint_t i) const { return m_size[i]; } - - GT_FUNCTION - uint_t perm(uint_t i) const { return m_perm[i]; } - - GT_FUNCTION - uint_t start(uint_t i) const { return m_start[i]; } + constexpr uint_t max() const { return m_sorted[0]; } + constexpr uint_t min() const { return m_sorted[2]; } + constexpr uint_t median() const { return m_sorted[1]; } + constexpr uint_t size(uint_t i) const { return m_size[i]; } + constexpr uint_t perm(uint_t i) const { return m_perm[i]; } + constexpr uint_t start(uint_t i) const { return m_start[i]; } }; array, 3>, 3> sizes; @@ -219,12 +208,12 @@ namespace gridtools { }; template - GT_FUNCTION shape_type const &shape(direction) const { + constexpr shape_type const &shape(direction) const { return sizes[I + 1][J + 1][K + 1]; } }; - GT_FUNCTION int thread_along_axis(int i, int j, int k, int axis) { + constexpr int thread_along_axis(int i, int j, int k, int axis) { assert(axis >= 0 && axis < 3); return axis == 0 ? i : axis == 1 ? j : k; } @@ -240,7 +229,7 @@ namespace gridtools { const uint_t i = blockIdx.x * apply_gpu_impl_::threads_per_block_x_t::value + threadIdx.x; const uint_t j = blockIdx.y * apply_gpu_impl_::threads_per_block_y_t::value + threadIdx.y; const uint_t k = blockIdx.z * apply_gpu_impl_::threads_per_block_z_t::value + threadIdx.z; - device::for_each([&](auto dir) { + for_each([&](auto dir) { if (predicate(dir)) { auto const &shape = conf.shape(dir); if (i < shape.max() && j < shape.median() && k < shape.min()) { @@ -299,7 +288,7 @@ namespace gridtools { (DataField0, Datafield1, DataField2, ...) */ template - void apply(DataFieldViews const &... data_field_views) const { + void apply(DataFieldViews const &...data_field_views) const { apply_gpu_impl_::loop_kernel<<>>( m_boundary_function, m_predicate, m_conf, data_field_views...); GT_CUDA_CHECK(cudaGetLastError()); diff --git a/include/gridtools/common/array.hpp b/include/gridtools/common/array.hpp index 429b910d85..790faaa6e7 100644 --- a/include/gridtools/common/array.hpp +++ b/include/gridtools/common/array.hpp @@ -12,8 +12,8 @@ @file @brief Implementation of an array class */ + #include -#include #include #include #include @@ -23,9 +23,7 @@ #include "../meta/list.hpp" #include "../meta/macros.hpp" #include "../meta/repeat.hpp" -#include "defs.hpp" #include "host_device.hpp" -#include "utility.hpp" namespace gridtools { @@ -61,32 +59,28 @@ namespace gridtools { using reverse_iterator = std::reverse_iterator; using const_reverse_iterator = std::reverse_iterator; - GT_FUNCTION constexpr T const *begin() const { return &m_array[0]; } + GT_FORCE_INLINE constexpr T const *begin() const { return &m_array[0]; } - GT_FUNCTION constexpr T *begin() { return &m_array[0]; } + GT_FORCE_INLINE constexpr T *begin() { return &m_array[0]; } - GT_FUNCTION constexpr T const *end() const { return &m_array[D]; } + GT_FORCE_INLINE constexpr T const *end() const { return &m_array[D]; } - GT_FUNCTION constexpr T *end() { return &m_array[D]; } + GT_FORCE_INLINE constexpr T *end() { return &m_array[D]; } - GT_FUNCTION constexpr const T *data() const noexcept { return m_array; } - GT_FUNCTION constexpr T *data() noexcept { return m_array; } + GT_FORCE_INLINE constexpr const T *data() const noexcept { return m_array; } + GT_FORCE_INLINE constexpr T *data() noexcept { return m_array; } - GT_FUNCTION GT_CONSTEXPR T const &operator[](size_t i) const { return m_array[i]; } + GT_FORCE_INLINE constexpr T const &operator[](size_t i) const { return m_array[i]; } - GT_FUNCTION T &operator[](size_t i) { - assert(i < D); - return m_array[i]; - } + GT_FORCE_INLINE constexpr T &operator[](size_t i) { return m_array[i]; } template - GT_FUNCTION array &operator=(A const &a) { - assert(a.size() == D); + GT_FORCE_INLINE constexpr array &operator=(A const &a) { std::copy(a.begin(), a.end(), m_array); return *this; } - GT_FUNCTION static constexpr size_t size() { return D; } + static GT_FORCE_INLINE constexpr size_t size() { return D; } }; template @@ -108,21 +102,21 @@ namespace gridtools { struct getter { template - static GT_FUNCTION GT_CONSTEXPR T &get(array &arr) noexcept { + static GT_FORCE_INLINE constexpr T &get(array &arr) noexcept { static_assert(I < D, "index is out of bounds"); return arr.m_array[I]; } template - static GT_FUNCTION GT_CONSTEXPR const T &get(const array &arr) noexcept { + static GT_FORCE_INLINE constexpr const T &get(const array &arr) noexcept { static_assert(I < D, "index is out of bounds"); return arr.m_array[I]; } template - static GT_FUNCTION GT_CONSTEXPR T &&get(array &&arr) noexcept { + static GT_FORCE_INLINE constexpr T &&get(array &&arr) noexcept { static_assert(I < D, "index is out of bounds"); - return wstd::move(arr.m_array[I]); + return std::move(arr.m_array[I]); } }; } // namespace array_impl_ @@ -138,7 +132,7 @@ namespace gridtools { // in case we need a constexpr version we need to implement a recursive one for c++11 template - GT_CONSTEXPR GT_FUNCTION bool operator==(array const &a, array const &b) { + GT_FORCE_INLINE constexpr bool operator==(array const &a, array const &b) { #ifndef __GNUC__ #pragma unroll #endif @@ -150,7 +144,7 @@ namespace gridtools { } template - GT_CONSTEXPR GT_FUNCTION bool operator!=(array const &a, array const &b) { + GT_FORCE_INLINE constexpr bool operator!=(array const &a, array const &b) { return !(a == b); } @@ -167,21 +161,21 @@ namespace gridtools { struct is_array_of, Value> : std::true_type {}; template - GT_FUNCTION T &get(array &arr) noexcept { + GT_FORCE_INLINE constexpr T &get(array &arr) noexcept { static_assert(I < D, "index is out of bounds"); return arr.m_array[I]; } template - GT_FUNCTION GT_CONSTEXPR const T &get(const array &arr) noexcept { + GT_FORCE_INLINE constexpr const T &get(const array &arr) noexcept { static_assert(I < D, "index is out of bounds"); return arr.m_array[I]; } template - GT_FUNCTION GT_CONSTEXPR T &&get(array &&arr) noexcept { + GT_FORCE_INLINE constexpr T &&get(array &&arr) noexcept { static_assert(I < D, "index is out of bounds"); - return wstd::move(get(arr)); + return std::move(get(arr)); } /** @} */ diff --git a/include/gridtools/common/atomic_host.hpp b/include/gridtools/common/atomic_host.hpp index 1d5a30c395..576e58d385 100644 --- a/include/gridtools/common/atomic_host.hpp +++ b/include/gridtools/common/atomic_host.hpp @@ -8,7 +8,11 @@ * SPDX-License-Identifier: BSD-3-Clause */ #pragma once + #include + +#include "host_device.hpp" + namespace gridtools { /** \ingroup common @@ -27,8 +31,7 @@ namespace gridtools { * @param val value added to var * @return the old value contained in var */ - GT_FUNCTION - static T atomic_add(T &var, const T val) { + static GT_FORCE_INLINE constexpr T atomic_add(T &var, const T val) { #if _OPENMP > 201106 T old; #pragma omp atomic capture @@ -54,8 +57,7 @@ namespace gridtools { * @param val value added to var * @return the old value contained in var */ - GT_FUNCTION - static T atomic_sub(T &var, const T val) { + static GT_FORCE_INLINE constexpr T atomic_sub(T &var, const T val) { #if _OPENMP > 201106 T old; #pragma omp atomic capture @@ -81,8 +83,7 @@ namespace gridtools { * @param val value inserted in variable var * @return the old value contained in var */ - GT_FUNCTION - static T atomic_exch(T &var, const T val) { + static GT_FORCE_INLINE constexpr T atomic_exch(T &var, const T val) { #if _OPENMP > 201106 T old; #pragma omp atomic capture @@ -108,8 +109,7 @@ namespace gridtools { * @param val value used in the min comparison * @return the old value contained in var */ - GT_FUNCTION - static T atomic_min(T &var, const T val) { + static GT_FORCE_INLINE constexpr T atomic_min(T &var, const T val) { T old; #pragma omp critical(min) { @@ -125,8 +125,7 @@ namespace gridtools { * @param val value used in the min comparison * @return the old value contained in var */ - GT_FUNCTION - static T atomic_max(T &var, const T val) { + static GT_FORCE_INLINE constexpr T atomic_max(T &var, const T val) { T old; #pragma omp critical(max) { diff --git a/include/gridtools/common/compose.hpp b/include/gridtools/common/compose.hpp index 4f9c127639..f87a0cbc8a 100644 --- a/include/gridtools/common/compose.hpp +++ b/include/gridtools/common/compose.hpp @@ -7,64 +7,23 @@ * Please, refer to the LICENSE file in the root directory. * SPDX-License-Identifier: BSD-3-Clause */ +#pragma once -#ifndef GT_TARGET_ITERATING -// DON'T USE #pragma once HERE!!! -#ifndef GT_COMMON_COMPOSE_HPP_ -#define GT_COMMON_COMPOSE_HPP_ - -#include "../meta.hpp" -#include "defs.hpp" -#include "host_device.hpp" -#include "utility.hpp" - -#define GT_FILENAME -#include GT_ITERATE_ON_TARGETS() -#undef GT_FILENAME - -#endif // GT_COMMON_COMPOSE_HPP_ -#else // GT_TARGET_ITERATING +#include namespace gridtools { - GT_TARGET_NAMESPACE { - namespace compose_impl_ { - template - struct composed_f; - - template - struct composed_f { - F m_f; - G m_g; - - template - GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE std::result_of_t)> - operator()(Args &&... args) const { - return m_f(m_g(wstd::forward(args)...)); - } - }; - - template - struct composed_f : composed_f> { - GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE composed_f(F f, Fs... fs) - : composed_f>{wstd::move(f), {wstd::move(fs)...}} {} - }; - } // namespace compose_impl_ - - /// Make function composition from provided functions - /// - /// compose(a, b, c)(x, y) <==> a(b(c(x, y))) - /// - template - GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE compose_impl_::composed_f...> compose( - Funs && ... funs) { - return {wstd::forward(funs)...}; - } + /// Make function composition from provided functions + /// + /// compose(a, b, c)(x, y) <==> a(b(c(x, y))) + /// + template + constexpr F &&compose(F &&f) { + return std::forward(f); + } - template - GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE Fun compose(Fun && fun) { - return fun; - } + template + constexpr auto compose(F &&f, Fs &&...fs) { + return [f = std::forward(f), fs = compose(std::forward(fs)...)]( + auto &&...args) { return f(fs(std::forward(args)...)); }; } } // namespace gridtools - -#endif // GT_TARGET_ITERATING diff --git a/include/gridtools/common/cuda_util.hpp b/include/gridtools/common/cuda_util.hpp index 30cc3b2574..907f55fbf4 100644 --- a/include/gridtools/common/cuda_util.hpp +++ b/include/gridtools/common/cuda_util.hpp @@ -60,14 +60,14 @@ namespace gridtools { return unique_cuda_ptr{ptr}; } - template ::value, int> = 0> + template , int> = 0> unique_cuda_ptr make_clone(T const &src) { unique_cuda_ptr res = cuda_malloc(); GT_CUDA_CHECK(cudaMemcpy(res.get(), &src, sizeof(T), cudaMemcpyHostToDevice)); return res; } - template ::value, int> = 0> + template , int> = 0> T from_clone(unique_cuda_ptr const &clone) { T res; GT_CUDA_CHECK(cudaMemcpy(&res, clone.get(), sizeof(T), cudaMemcpyDeviceToHost)); diff --git a/include/gridtools/common/defs.hpp b/include/gridtools/common/defs.hpp index d96f34806b..957f9b1465 100644 --- a/include/gridtools/common/defs.hpp +++ b/include/gridtools/common/defs.hpp @@ -26,8 +26,6 @@ namespace gridtools { #endif #endif -#define GT_CONSTEXPR constexpr - #ifdef __cpp_consteval #define GT_CONSTEVAL consteval #else diff --git a/include/gridtools/common/enable_maker.hpp b/include/gridtools/common/enable_maker.hpp index f3d030b56e..50908b6800 100644 --- a/include/gridtools/common/enable_maker.hpp +++ b/include/gridtools/common/enable_maker.hpp @@ -21,7 +21,7 @@ namespace gridtools { template