/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/utility/unary_element_function.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/utility/unary_element_function.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/utility/unary_element_function.hpp Source File
unary_element_function.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck_tile/core.hpp"
7 
8 namespace ck_tile {
9 
10 template <typename F, typename... Fs>
11 struct composes : private composes<F>
12 {
13  template <typename FirstArg, typename... RestArgs>
14  CK_TILE_HOST_DEVICE constexpr explicit composes(FirstArg&& firstArg, RestArgs&&... restArgs)
15  : composes<F>(std::forward<FirstArg>(firstArg)), inner_(std::forward<RestArgs>(restArgs)...)
16  {
17  }
18 
19  template <typename Arg>
20  CK_TILE_HOST_DEVICE constexpr auto operator()(Arg&& arg) const
21  {
22  return static_cast<const composes<F>&>(*this)(inner_(std::forward<Arg>(arg)));
23  }
24 
25  private:
26  composes<Fs...> inner_;
27 };
28 
29 template <typename F>
30 struct composes<F>
31 {
32  static_assert(!std::is_reference_v<F>);
33 
34  template <typename Arg, typename = std::enable_if_t<std::is_constructible_v<F, Arg>>>
35  CK_TILE_HOST_DEVICE constexpr explicit composes(Arg&& arg) : f_(std::forward<Arg>(arg))
36  {
37  }
38 
39  template <typename Arg,
40  typename = std::enable_if_t<std::is_invocable_v<std::add_const_t<F>&, Arg>>>
41  CK_TILE_HOST_DEVICE constexpr auto operator()(Arg&& arg) const
42  {
43  return f_(std::forward<Arg>(arg));
44  }
45 
46  private:
47  F f_;
48 };
49 
51 template <typename... Ts>
52 __host__ __device__ composes(Ts&&...) -> composes<remove_cvref_t<Ts>...>;
53 
54 template <typename SaturateType>
55 struct saturates
56 {
57  // NOTE: this function does not return SaturateType value
58  // it is user's responsiblity to do further cast or not
59  template <typename AccType>
60  CK_TILE_HOST_DEVICE constexpr auto
61  operator()(const AccType& a_) const -> std::enable_if_t<std::is_arithmetic_v<AccType>, AccType>
62  {
63  return clamp(a_,
64  type_convert<AccType>(numeric<SaturateType>::lowest()),
65  type_convert<AccType>(numeric<SaturateType>::max()));
66  }
67 };
68 
69 } // namespace ck_tile
#define CK_TILE_HOST_DEVICE
Definition: config.hpp:42
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_HOST_DEVICE T clamp(const T &x, const T &lowerbound, const T &upperbound)
Definition: math.hpp:259
__host__ __device__ composes(Ts &&...) -> composes< remove_cvref_t< Ts >... >
FIXME: create macro to replace 'host device' and nothing more.
Definition: unary_element_function.hpp:31
constexpr CK_TILE_HOST_DEVICE composes(Arg &&arg)
Definition: unary_element_function.hpp:35
constexpr CK_TILE_HOST_DEVICE auto operator()(Arg &&arg) const
Definition: unary_element_function.hpp:41
Definition: unary_element_function.hpp:12
constexpr CK_TILE_HOST_DEVICE auto operator()(Arg &&arg) const
Definition: unary_element_function.hpp:20
constexpr CK_TILE_HOST_DEVICE composes(FirstArg &&firstArg, RestArgs &&... restArgs)
Definition: unary_element_function.hpp:14
Definition: numeric.hpp:18
Definition: unary_element_function.hpp:56
constexpr CK_TILE_HOST_DEVICE auto operator()(const AccType &a_) const -> std::enable_if_t< std::is_arithmetic_v< AccType >, AccType >
Definition: unary_element_function.hpp:61