math.hpp File Reference

math.hpp File Reference#

Composable Kernel: math.hpp File Reference
Classes | Namespaces | Macros | Functions | Variables
math.hpp File Reference
#include "ck_tile/core/config.hpp"
#include "ck_tile/core/numeric/integer.hpp"
#include "ck_tile/core/numeric/integral_constant.hpp"
#include "ck_tile/core/utility/bit_cast.hpp"
#include <type_traits>
#include <stdint.h>
#include <cmath>
#include "ck_tile/core/numeric/half.hpp"
#include "ck_tile/core/numeric/bfloat16.hpp"
#include "ck_tile/core/numeric/float8.hpp"
#include "ck_tile/core/numeric/type_convert.hpp"

Go to the source code of this file.

Classes

struct  ck_tile::scales_c< Scale, lhs >
struct  ck_tile::scales< Scale >
struct  ck_tile::plus< Left, Right >
struct  ck_tile::plus< void, void >
struct  ck_tile::minus< Left, Right >
struct  ck_tile::minus< void, void >
struct  ck_tile::multiplies< Left, Right >
struct  ck_tile::multiplies< void, void >
struct  ck_tile::maximize< T >
struct  ck_tile::minimize< T >
struct  ck_tile::integer_divide_ceiler< T >
struct  ck_tile::equal< Left, Right >
struct  ck_tile::equal< void, void >
struct  ck_tile::equal< float, float >
struct  ck_tile::equal< double, double >
struct  ck_tile::less< Left, Right >
struct  ck_tile::less< void, void >
struct  ck_tile::less_equal< Left, Right >
struct  ck_tile::less_equal< void, void >
struct  ck_tile::less_equal< float, float >
struct  ck_tile::less_equal< double, double >
struct  ck_tile::log2e< double >
struct  ck_tile::log2e< float >

Namespaces

namespace  ck_tile

Macros

#define C_LOG2E   1.44269504088896340736

Functions

template<typename Scale>
__host__ __device__ ck_tile::scales (Scale) -> scales< Scale >
 FIXME: create macro to replace 'host device' and nothing more.
__host__ __device__ ck_tile::plus () -> plus< void, void >
 FIXME: create macro to replace 'host device' and nothing more.
__host__ __device__ ck_tile::minus () -> minus< void, void >
 FIXME: create macro to replace 'host device' and nothing more.
__host__ __device__ ck_tile::multiplies () -> multiplies< void, void >
 FIXME: create macro to replace 'host device' and nothing more.
template<typename X, typename Y>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::integer_divide_floor (X x, Y y)
template<typename X, typename Y>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::integer_divide_ceil (X x, Y y)
template<typename X, typename Y>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::integer_least_multiple (X x, Y y)
template<typename T>
CK_TILE_HOST_DEVICE constexpr T ck_tile::max (T x)
template<typename T>
CK_TILE_HOST constexpr T ck_tile::max (T x, T y)
template<typename T>
CK_TILE_DEVICE constexpr T ck_tile::max (T x, T y)
template<>
CK_TILE_DEVICE constexpr float ck_tile::max (float x, float y)
template<>
CK_TILE_DEVICE constexpr double ck_tile::max (double x, double y)
template<index_t X>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::max (number< X >, index_t y)
template<index_t Y>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::max (index_t x, number< Y >)
template<typename X, typename... Ys>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::max (X x, Ys... ys)
template<typename T>
CK_TILE_HOST_DEVICE constexpr T ck_tile::min (T x)
template<typename T>
CK_TILE_HOST constexpr T ck_tile::min (T x, T y)
template<typename T>
CK_TILE_DEVICE constexpr T ck_tile::min (T x, T y)
template<>
CK_TILE_DEVICE constexpr float ck_tile::min (float x, float y)
template<>
CK_TILE_DEVICE constexpr double ck_tile::min (double x, double y)
template<index_t X>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::min (number< X >, index_t y)
template<index_t Y>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::min (index_t x, number< Y >)
template<typename X, typename... Ys>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::min (X x, Ys... ys)
template<typename T>
CK_TILE_HOST_DEVICE constexpr T ck_tile::clamp (const T &x, const T &lowerbound, const T &upperbound)
CK_TILE_HOST int ck_tile::clz (uint32_t x)
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::gcd (index_t x, index_t y)
template<index_t X, index_t Y>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::gcd (number< X >, number< Y >)
template<typename X, typename... Ys, typename std::enable_if< sizeof...(Ys) > = 2, bool, ::type = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::gcd (X x, Ys... ys)
template<typename X, typename Y>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::lcm (X x, Y y)
template<typename X, typename... Ys, typename std::enable_if< sizeof...(Ys) > = 2, bool, ::type = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::lcm (X x, Ys... ys)
__host__ __device__ ck_tile::equal () -> equal< void, void >
 FIXME: create macro to replace 'host device' and nothing more.
__host__ __device__ ck_tile::less () -> less< void, void >
 FIXME: create macro to replace 'host device' and nothing more.
__host__ __device__ ck_tile::less_equal () -> less_equal< void, void >
 FIXME: create macro to replace 'host device' and nothing more.
CK_TILE_HOST_DEVICE constexpr int32_t ck_tile::next_power_of_two (int32_t x)
template<index_t X>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::next_power_of_two ()
template<index_t X>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::next_power_of_two (number< X >)
CK_TILE_HOST_DEVICE constexpr int32_t ck_tile::integer_log2_floor (int32_t x)
CK_TILE_HOST_DEVICE constexpr bool ck_tile::is_power_of_two_integer (int32_t x)
CK_TILE_DEVICE float ck_tile::exp2 (float x)
CK_TILE_DEVICE uint16_t ck_tile::sad_u16 (uint16_t x, uint16_t y, uint16_t acc)
CK_TILE_DEVICE uint32_t ck_tile::sad_u32 (uint32_t x, uint32_t y, uint32_t acc)
CK_TILE_HOST float ck_tile::abs (float x)
CK_TILE_HOST double ck_tile::abs (double x)
CK_TILE_HOST int8_t ck_tile::abs (int8_t x)
CK_TILE_HOST int32_t ck_tile::abs (int32_t x)
CK_TILE_HOST fp16_t ck_tile::abs (fp16_t x)
CK_TILE_HOST bool ck_tile::isnan (float x)
CK_TILE_HOST bool ck_tile::isnan (double x)
CK_TILE_HOST bool ck_tile::isnan (int8_t x)
CK_TILE_HOST bool ck_tile::isnan (int32_t x)
CK_TILE_HOST bool ck_tile::isnan (fp16_t x)
CK_TILE_HOST fp16_t ck_tile::sqrt (fp16_t x)
CK_TILE_HOST float ck_tile::sqrt (float x)
CK_TILE_HOST double ck_tile::sqrt (double x)
template<typename T>
CK_TILE_HOST T ck_tile::tanh (T x)
template<>
CK_TILE_HOST float ck_tile::tanh< float > (float x)
template<>
CK_TILE_HOST double ck_tile::tanh< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::acos (T x)
template<>
CK_TILE_HOST float ck_tile::acos< float > (float x)
template<>
CK_TILE_HOST double ck_tile::acos< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::neg (T x)
template<>
CK_TILE_HOST float ck_tile::neg< float > (float x)
template<>
CK_TILE_HOST double ck_tile::neg< double > (double x)
template<>
CK_TILE_HOST int32_t ck_tile::neg< int32_t > (int32_t x)
template<>
CK_TILE_HOST int8_t ck_tile::neg< int8_t > (int8_t x)
template<typename T>
CK_TILE_HOST T ck_tile::atan (T x)
template<>
CK_TILE_HOST float ck_tile::atan< float > (float x)
template<>
CK_TILE_HOST double ck_tile::atan< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::sin (T x)
template<>
CK_TILE_HOST float ck_tile::sin< float > (float x)
template<>
CK_TILE_HOST double ck_tile::sin< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::asin (T x)
template<>
CK_TILE_HOST float ck_tile::asin< float > (float x)
template<>
CK_TILE_HOST double ck_tile::asin< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::asinh (T x)
template<>
CK_TILE_HOST float ck_tile::asinh< float > (float x)
template<>
CK_TILE_HOST double ck_tile::asinh< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::cos (T x)
template<>
CK_TILE_HOST float ck_tile::cos< float > (float x)
template<>
CK_TILE_HOST double ck_tile::cos< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::acosh (T x)
template<>
CK_TILE_HOST float ck_tile::acosh< float > (float x)
template<>
CK_TILE_HOST double ck_tile::acosh< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::tan (T x)
template<>
CK_TILE_HOST float ck_tile::tan< float > (float x)
template<>
CK_TILE_HOST double ck_tile::tan< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::atanh (T x)
template<>
CK_TILE_HOST float ck_tile::atanh< float > (float x)
template<>
CK_TILE_HOST double ck_tile::atanh< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::sinh (T x)
template<>
CK_TILE_HOST float ck_tile::sinh< float > (float x)
template<>
CK_TILE_HOST double ck_tile::sinh< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::ceil (T x)
template<>
CK_TILE_HOST float ck_tile::ceil< float > (float x)
template<>
CK_TILE_HOST double ck_tile::ceil< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::cosh (T x)
template<>
CK_TILE_HOST float ck_tile::cosh< float > (float x)
template<>
CK_TILE_HOST double ck_tile::cosh< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::floor (T x)
template<>
CK_TILE_HOST float ck_tile::floor< float > (float x)
template<>
CK_TILE_HOST double ck_tile::floor< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::rcp (T x)
template<typename T>
CK_TILE_HOST T ck_tile::exp (T x)
template<>
CK_TILE_HOST float ck_tile::exp< float > (float x)
template<>
CK_TILE_HOST double ck_tile::exp< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::log (T x)
template<>
CK_TILE_HOST float ck_tile::log< float > (float x)
template<>
CK_TILE_HOST double ck_tile::log< double > (double x)
template<typename T>
CK_TILE_HOST T ck_tile::pow (T x, T gamma)
template<>
CK_TILE_HOST float ck_tile::pow< float > (float x, float gamma)
template<>
CK_TILE_HOST double ck_tile::pow< double > (double x, double gamma)
template<typename T>
CK_TILE_HOST T ck_tile::expm1 (T x)
template<>
CK_TILE_HOST float ck_tile::expm1< float > (float x)
template<>
CK_TILE_HOST double ck_tile::expm1< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::tanh (T x)
template<>
CK_TILE_DEVICE float ck_tile::tanh< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::tanh< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::acos (T x)
template<>
CK_TILE_DEVICE float ck_tile::acos< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::acos< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::neg (T x)
template<>
CK_TILE_DEVICE float ck_tile::neg< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::neg< double > (double x)
template<>
CK_TILE_DEVICE int32_t ck_tile::neg< int32_t > (int32_t x)
template<>
CK_TILE_DEVICE int8_t ck_tile::neg< int8_t > (int8_t x)
template<>
CK_TILE_DEVICE fp16_t ck_tile::neg< fp16_t > (fp16_t x)
template<typename T>
CK_TILE_DEVICE T ck_tile::atan (T x)
template<>
CK_TILE_DEVICE float ck_tile::atan< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::atan< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::sin (T x)
template<>
CK_TILE_DEVICE float ck_tile::sin< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::sin< double > (double x)
template<>
CK_TILE_DEVICE fp16_t ck_tile::sin< fp16_t > (fp16_t x)
template<typename T>
CK_TILE_DEVICE T ck_tile::asin (T x)
template<>
CK_TILE_DEVICE float ck_tile::asin< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::asin< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::asinh (T x)
template<>
CK_TILE_DEVICE float ck_tile::asinh< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::asinh< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::acosh (T x)
template<>
CK_TILE_DEVICE float ck_tile::acosh< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::acosh< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::tan (T x)
template<>
CK_TILE_DEVICE float ck_tile::tan< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::tan< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::atanh (T x)
template<>
CK_TILE_DEVICE float ck_tile::atanh< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::atanh< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::sinh (T x)
template<>
CK_TILE_DEVICE float ck_tile::sinh< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::sinh< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::ceil (T x)
template<>
CK_TILE_DEVICE float ck_tile::ceil< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::ceil< double > (double x)
template<>
CK_TILE_DEVICE fp16_t ck_tile::ceil< fp16_t > (fp16_t x)
template<typename T>
CK_TILE_DEVICE T ck_tile::cosh (T x)
template<>
CK_TILE_DEVICE float ck_tile::cosh< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::cosh< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::floor (T x)
template<>
CK_TILE_DEVICE float ck_tile::floor< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::floor< double > (double x)
template<>
CK_TILE_DEVICE fp16_t ck_tile::floor< fp16_t > (fp16_t x)
template<typename T>
CK_TILE_DEVICE T ck_tile::rcp (T x)
template<typename T>
CK_TILE_DEVICE T ck_tile::exp (T x)
template<>
CK_TILE_DEVICE fp16_t ck_tile::exp< fp16_t > (fp16_t x)
template<>
CK_TILE_DEVICE float ck_tile::exp< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::exp< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::tanh_fast (T x)
template<>
CK_TILE_DEVICE float ck_tile::tanh_fast< float > (float x)
template<typename T>
CK_TILE_DEVICE T ck_tile::log (T x)
template<>
CK_TILE_DEVICE fp16_t ck_tile::log< fp16_t > (fp16_t x)
template<>
CK_TILE_DEVICE float ck_tile::log< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::log< double > (double x)
template<typename T>
CK_TILE_DEVICE T ck_tile::pow (T x, T gamma)
template<>
CK_TILE_DEVICE float ck_tile::pow< float > (float x, float gamma)
template<>
CK_TILE_DEVICE double ck_tile::pow< double > (double x, double gamma)
template<typename T>
CK_TILE_DEVICE T ck_tile::expm1 (T x)
template<>
CK_TILE_DEVICE float ck_tile::expm1< float > (float x)
template<>
CK_TILE_DEVICE double ck_tile::expm1< double > (double x)

Variables

template<typename T = double>
constexpr T ck_tile::log2e_v = log2e<T>::value
template<typename T = double>
constexpr T ck_tile::log2e_rcp_v = 1. / log2e<T>::value

Macro Definition Documentation

◆ C_LOG2E

#define C_LOG2E   1.44269504088896340736