Constexpr

examples/boundaries/boundaries.cpp

@@ -64,11 +64,9 @@ template <typename T>
 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.

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();

examples/stencil_type_erasure/interpolate_stencil.cpp

@@ -48,6 +48,6 @@ namespace {
 std::function<void(inputs, outputs)> 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);
     };
 }

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 <sign I, sign J, sign K>
-                GT_FUNCTION static constexpr uint_t direction_index(direction<I, J, K>) {
+                static constexpr uint_t direction_index(direction<I, J, K>) {
                     // 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<array<array<shape_type, 3>, 3>, 3> sizes;
@@ -219,12 +208,12 @@ namespace gridtools {
                 };
 
                 template <sign I, sign J, sign K>
-                GT_FUNCTION shape_type const &shape(direction<I, J, K>) const {
+                constexpr shape_type const &shape(direction<I, J, K>) 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<directions_t>([&](auto dir) {
+                for_each<directions_t>([&](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 <typename... DataFieldViews>
-            void apply(DataFieldViews const &... data_field_views) const {
+            void apply(DataFieldViews const &...data_field_views) const {
                 apply_gpu_impl_::loop_kernel<<<m_conf.kernel_grid_size(), m_conf.kernel_thread_block_size()>>>(
                     m_boundary_function, m_predicate, m_conf, data_field_views...);
                 GT_CUDA_CHECK(cudaGetLastError());

include/gridtools/common/array.hpp

@@ -12,8 +12,8 @@
 @file
 @brief Implementation of an array class
 */
+
 #include <algorithm>
-#include <cassert>
 #include <iterator>
 #include <tuple>
 #include <type_traits>
@@ -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<iterator>;
         using const_reverse_iterator = std::reverse_iterator<const_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 <typename A>
-        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 <class T, class... Ts>
@@ -108,21 +102,21 @@ namespace gridtools {
 
         struct getter {
             template <size_t I, typename T, size_t D>
-            static GT_FUNCTION GT_CONSTEXPR T &get(array<T, D> &arr) noexcept {
+            static GT_FORCE_INLINE constexpr T &get(array<T, D> &arr) noexcept {
                 static_assert(I < D, "index is out of bounds");
                 return arr.m_array[I];
             }
 
             template <size_t I, typename T, size_t D>
-            static GT_FUNCTION GT_CONSTEXPR const T &get(const array<T, D> &arr) noexcept {
+            static GT_FORCE_INLINE constexpr const T &get(const array<T, D> &arr) noexcept {
                 static_assert(I < D, "index is out of bounds");
                 return arr.m_array[I];
             }
 
             template <size_t I, typename T, size_t D>
-            static GT_FUNCTION GT_CONSTEXPR T &&get(array<T, D> &&arr) noexcept {
+            static GT_FORCE_INLINE constexpr T &&get(array<T, D> &&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 <typename T, typename U, size_t D>
-    GT_CONSTEXPR GT_FUNCTION bool operator==(array<T, D> const &a, array<U, D> const &b) {
+    GT_FORCE_INLINE constexpr bool operator==(array<T, D> const &a, array<U, D> const &b) {
 #ifndef __GNUC__
 #pragma unroll
 #endif
@@ -150,7 +144,7 @@ namespace gridtools {
     }
 
     template <typename T, typename U, size_t D>
-    GT_CONSTEXPR GT_FUNCTION bool operator!=(array<T, D> const &a, array<U, D> const &b) {
+    GT_FORCE_INLINE constexpr bool operator!=(array<T, D> const &a, array<U, D> const &b) {
         return !(a == b);
     }
 
@@ -167,21 +161,21 @@ namespace gridtools {
     struct is_array_of<array<Value, D>, Value> : std::true_type {};
 
     template <size_t I, typename T, size_t D>
-    GT_FUNCTION T &get(array<T, D> &arr) noexcept {
+    GT_FORCE_INLINE constexpr T &get(array<T, D> &arr) noexcept {
         static_assert(I < D, "index is out of bounds");
         return arr.m_array[I];
     }
 
     template <size_t I, typename T, size_t D>
-    GT_FUNCTION GT_CONSTEXPR const T &get(const array<T, D> &arr) noexcept {
+    GT_FORCE_INLINE constexpr const T &get(const array<T, D> &arr) noexcept {
         static_assert(I < D, "index is out of bounds");
         return arr.m_array[I];
     }
 
     template <size_t I, typename T, size_t D>
-    GT_FUNCTION GT_CONSTEXPR T &&get(array<T, D> &&arr) noexcept {
+    GT_FORCE_INLINE constexpr T &&get(array<T, D> &&arr) noexcept {
         static_assert(I < D, "index is out of bounds");
-        return wstd::move(get<I>(arr));
+        return std::move(get<I>(arr));
     }
 
     /** @} */

include/gridtools/common/atomic_host.hpp

@@ -8,7 +8,11 @@
  * SPDX-License-Identifier: BSD-3-Clause
  */
 #pragma once
+
 #include <algorithm>
+
+#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)
             {

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 <gridtools/common/compose.hpp>
-#include GT_ITERATE_ON_TARGETS()
-#undef GT_FILENAME
-
-#endif // GT_COMMON_COMPOSE_HPP_
-#else  // GT_TARGET_ITERATING
+#include <utility>
 
 namespace gridtools {
-    GT_TARGET_NAMESPACE {
-        namespace compose_impl_ {
-            template <class... Funs>
-            struct composed_f;
-
-            template <class F, class G>
-            struct composed_f<F, G> {
-                F m_f;
-                G m_g;
-
-                template <class... Args>
-                GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE std::result_of_t<F(std::result_of_t<G(Args &&...)>)>
-                operator()(Args &&... args) const {
-                    return m_f(m_g(wstd::forward<Args>(args)...));
-                }
-            };
-
-            template <class F, class... Fs>
-            struct composed_f<F, Fs...> : composed_f<F, composed_f<Fs...>> {
-                GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE composed_f(F f, Fs... fs)
-                    : composed_f<F, composed_f<Fs...>>{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 <class... Funs>
-        GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE compose_impl_::composed_f<std::decay_t<Funs>...> compose(
-            Funs && ... funs) {
-            return {wstd::forward<Funs>(funs)...};
-        }
+    /// Make function composition from provided functions
+    ///
+    /// compose(a, b, c)(x, y) <==> a(b(c(x, y)))
+    ///
+    template <class F>
+    constexpr F &&compose(F &&f) {
+        return std::forward<F>(f);
+    }
 
-        template <class Fun>
-        GT_TARGET_CONSTEXPR GT_TARGET GT_FORCE_INLINE Fun compose(Fun && fun) {
-            return fun;
-        }
+    template <class F, class... Fs>
+    constexpr auto compose(F &&f, Fs &&...fs) {
+        return [f = std::forward<F>(f), fs = compose(std::forward<Fs>(fs)...)](
+                   auto &&...args) { return f(fs(std::forward<decltype(args)>(args)...)); };
     }
 } // namespace gridtools
-
-#endif // GT_TARGET_ITERATING

include/gridtools/common/cuda_util.hpp

@@ -60,14 +60,14 @@ namespace gridtools {
             return unique_cuda_ptr<T>{ptr};
         }
 
-        template <class T, std::enable_if_t<std::is_trivially_copyable<T>::value, int> = 0>
+        template <class T, std::enable_if_t<std::is_trivially_copy_constructible_v<T>, int> = 0>
         unique_cuda_ptr<T> make_clone(T const &src) {
             unique_cuda_ptr<T> res = cuda_malloc<T>();
             GT_CUDA_CHECK(cudaMemcpy(res.get(), &src, sizeof(T), cudaMemcpyHostToDevice));
             return res;
         }
 
-        template <class T, std::enable_if_t<std::is_trivially_copyable<T>::value, int> = 0>
+        template <class T, std::enable_if_t<std::is_trivially_copy_constructible_v<T>, int> = 0>
         T from_clone(unique_cuda_ptr<T> const &clone) {
             T res;
             GT_CUDA_CHECK(cudaMemcpy(&res, clone.get(), sizeof(T), cudaMemcpyDeviceToHost));