// SPDX-License-Identifier: MIT // Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved. #ifndef CK_STATICALLY_INDEXED_ARRAY_MULTI_INDEX_HPP #define CK_STATICALLY_INDEXED_ARRAY_MULTI_INDEX_HPP #include "common_header.hpp" namespace ck { template using MultiIndex = StaticallyIndexedArray; template __host__ __device__ constexpr auto make_multi_index(Xs&&... xs) { return make_statically_indexed_array(index_t{xs}...); } template __host__ __device__ constexpr auto make_zero_multi_index() { return unpack([](auto... xs) { return make_multi_index(xs...); }, typename uniform_sequence_gen::type{}); } template __host__ __device__ constexpr auto to_multi_index(const T& x) { return unpack([](auto... ys) { return make_multi_index(ys...); }, x); } // Here should use MultiIndex, instead of Tuple, although the former // is the alias of the latter. This is because compiler cannot infer the NSize if // using MultiIndex // TODO: how to fix this? template < typename... Ys, typename X, enable_if_t::value && !std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator+=(Tuple& y, const X& x) { static_assert(X::Size() == sizeof...(Ys), "wrong! size not the same"); constexpr index_t NSize = sizeof...(Ys); static_for<0, NSize, 1>{}([&](auto i) { y(i) += x[i]; }); return y; } template < typename... Ys, typename X, enable_if_t::value && !std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator-=(Tuple& y, const X& x) { static_assert(X::Size() == sizeof...(Ys), "wrong! size not the same"); constexpr index_t NSize = sizeof...(Ys); static_for<0, NSize, 1>{}([&](auto i) { y(i) -= x[i]; }); return y; } template < typename... Xs, typename Y, enable_if_t::value && !std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator+(const Tuple& x, const Y& y) { static_assert(Y::Size() == sizeof...(Xs), "wrong! size not the same"); constexpr index_t NSize = sizeof...(Xs); Tuple r; static_for<0, NSize, 1>{}([&](auto i) { r(i) = x[i] + y[i]; }); return r; } template < typename... Xs, typename Y, enable_if_t::value && !std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator-(const Tuple& x, const Y& y) { static_assert(Y::Size() == sizeof...(Xs), "wrong! size not the same"); constexpr index_t NSize = sizeof...(Xs); Tuple r; static_for<0, NSize, 1>{}([&](auto i) { r(i) = x[i] - y[i]; }); return r; } template < typename... Xs, typename Y, enable_if_t::value && !std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator*(const Tuple& x, const Y& y) { static_assert(Y::Size() == sizeof...(Xs), "wrong! size not the same"); constexpr index_t NSize = sizeof...(Xs); Tuple r; static_for<0, NSize, 1>{}([&](auto i) { r(i) = x[i] * y[i]; }); return r; } // MultiIndex = scalar * MultiIndex template ::value || std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator*(Y a, const Tuple& x) { constexpr index_t NSize = sizeof...(Xs); Tuple r; static_for<0, NSize, 1>{}([&](auto i) { r(i) = a * x[i]; }); return r; } // MultiIndex = MultiIndex * scalar template ::value || std::is_floating_point::value, bool> = false> __host__ __device__ constexpr auto operator*(const Tuple& x, Y a) { return a * x; } namespace mathext { template __host__ __device__ constexpr auto exp(const Tuple& x) { constexpr index_t NSize = sizeof...(Xs); Tuple r; static_for<0, NSize, 1>{}([&](auto i) { r(i) = math::exp(x[i]); }); return r; } template __host__ __device__ constexpr auto max(const Tuple& x, const Y& y) { static_assert(Y::Size() == sizeof...(Xs), "wrong! size not the same"); constexpr index_t NSize = sizeof...(Xs); Tuple r; static_for<0, NSize, 1>{}([&](auto i) { r(i) = math::max(x[i], y[i]); }); return r; } } // namespace mathext template __host__ __device__ void print_multi_index(const Tuple& x) { printf("{"); printf("MultiIndex, "); printf("size %d,", index_t{sizeof...(Xs)}); static_for<0, sizeof...(Xs), 1>{}( [&](auto i) { printf("%d ", static_cast(x.At(i))); }); printf("}"); } } // namespace ck #endif