// SPDX-License-Identifier: MIT // Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved. #pragma once #ifndef __HIP_DEVICE_COMPILE__ #include #endif #include "ck/utility/data_type.hpp" #include "ck/utility/type.hpp" #include "ck/utility/type_convert.hpp" namespace ck { namespace math { #if CK_WORKAROUND_SWDEV_383542 extern "C" __device__ float __ocml_native_recip_f32(float); #endif // math functions for the host, some are implemented by calling C++ std functions static inline __host__ float abs(float x) { return std::abs(x); }; static inline __host__ double abs(double x) { return std::abs(x); }; static inline __host__ int8_t abs(int8_t x) { int8_t sgn = x >> (8 - 1); return (x ^ sgn) - sgn; }; static inline __host__ int32_t abs(int32_t x) { int32_t sgn = x >> (32 - 1); return (x ^ sgn) - sgn; }; static inline __host__ half_t abs(half_t x) { uint16_t xx = ck::bit_cast(x); uint16_t abs_xx = xx & 0x7fff; half_t abs_x = ck::bit_cast(abs_xx); return abs_x; }; #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 static inline __host__ int4_t abs(int4_t x) { int4_t sgn = x >> (4 - 1); return (x ^ sgn) - sgn; } #endif static inline __host__ bool isnan(float x) { return std::isnan(x); }; static inline __host__ bool isnan(double x) { return std::isnan(x); }; static inline __host__ bool isnan(int8_t x) { (void)x; return false; }; static inline __host__ bool isnan(int32_t x) { (void)x; return false; }; static inline __host__ bool isnan(half_t x) { uint16_t xx = ck::bit_cast(x); return (xx & 0x7FFF) > 0x7C00; }; static inline __host__ bool isnan(f8_t x) { return ck::fp8_is_nan(x); }; #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 static inline __host__ bool isnan(int4_t x) { (void)x; return false; }; #endif static inline __host__ half_t sqrt(half_t x) { return static_cast(std::sqrt(static_cast(x))); }; static inline __host__ float sqrt(float x) { return std::sqrt(x); }; static inline __host__ double sqrt(double x) { return std::sqrt(x); }; template inline __host__ T tanh(T x) { return ck::type_convert(std::tanhf(ck::type_convert(x))); }; template <> inline __host__ float tanh(float x) { return std::tanhf(x); }; template <> inline __host__ double tanh(double x) { return std::tanh(x); }; template inline __host__ T acos(T x) { return ck::type_convert(std::acosf(ck::type_convert(x))); }; template <> inline __host__ float acos(float x) { return std::acosf(x); }; template <> inline __host__ double acos(double x) { return std::acos(x); }; template inline __host__ T neg(T x) { return ck::type_convert(-(ck::type_convert(x))); }; template <> inline __host__ float neg(float x) { return -x; }; template <> inline __host__ double neg(double x) { return -x; }; template <> inline __host__ int32_t neg(int32_t x) { return -x; }; template <> inline __host__ int8_t neg(int8_t x) { return -x; }; template inline __host__ T atan(T x) { return ck::type_convert(std::atanf(ck::type_convert(x))); }; template <> inline __host__ float atan(float x) { return std::atanf(x); }; template <> inline __host__ double atan(double x) { return std::atan(x); }; template inline __host__ T sin(T x) { return ck::type_convert(std::sinf(ck::type_convert(x))); }; template <> inline __host__ float sin(float x) { return std::sinf(x); }; template <> inline __host__ double sin(double x) { return std::sin(x); }; template inline __host__ T asin(T x) { return ck::type_convert(std::asinf(ck::type_convert(x))); }; template <> inline __host__ float asin(float x) { return std::asinf(x); }; template <> inline __host__ double asin(double x) { return std::asin(x); }; template inline __host__ T asinh(T x) { return ck::type_convert(std::asinhf(ck::type_convert(x))); }; template <> inline __host__ float asinh(float x) { return std::asinhf(x); }; template <> inline __host__ double asinh(double x) { return std::asinh(x); }; template inline __host__ T cos(T x) { return ck::type_convert(std::cosf(ck::type_convert(x))); }; template <> inline __host__ float cos(float x) { return std::cosf(x); }; template <> inline __host__ double cos(double x) { return std::cos(x); }; template inline __host__ T acosh(T x) { return ck::type_convert(std::acoshf(ck::type_convert(x))); }; template <> inline __host__ float acosh(float x) { return std::acoshf(x); }; template <> inline __host__ double acosh(double x) { return std::acosh(x); }; template inline __host__ T tan(T x) { return ck::type_convert(std::tanf(ck::type_convert(x))); }; template <> inline __host__ float tan(float x) { return std::tanf(x); }; template <> inline __host__ double tan(double x) { return std::tan(x); }; template inline __host__ T atanh(T x) { return ck::type_convert(std::atanhf(ck::type_convert(x))); }; template <> inline __host__ float atanh(float x) { return std::atanhf(x); }; template <> inline __host__ double atanh(double x) { return std::atanh(x); }; template inline __host__ T sinh(T x) { return ck::type_convert(std::sinhf(ck::type_convert(x))); }; template <> inline __host__ float sinh(float x) { return std::sinhf(x); }; template <> inline __host__ double sinh(double x) { return std::sinh(x); }; template inline __host__ T ceil(T x) { return ck::type_convert(std::ceilf(ck::type_convert(x))); }; template <> inline __host__ float ceil(float x) { return std::ceilf(x); }; template <> inline __host__ double ceil(double x) { return std::ceil(x); }; template inline __host__ T cosh(T x) { return ck::type_convert(std::coshf(ck::type_convert(x))); }; template <> inline __host__ float cosh(float x) { return std::coshf(x); }; template <> inline __host__ double cosh(double x) { return std::cosh(x); }; template inline __host__ T floor(T x) { return ck::type_convert(std::floorf(ck::type_convert(x))); }; template <> inline __host__ float floor(float x) { return std::floorf(x); }; template <> inline __host__ double floor(double x) { return std::floor(x); }; template inline __host__ T rcp(T x) { return ck::type_convert(1.f / ck::type_convert(x)); }; template inline __host__ T exp(T x) { return ck::type_convert(std::expf(ck::type_convert(x))); } template <> inline __host__ float exp(float x) { return std::expf(x); } template <> inline __host__ double exp(double x) { return std::exp(x); } template inline __host__ T log(T x) { return ck::type_convert(std::logf(ck::type_convert(x))); } template <> inline __host__ float log(float x) { return std::logf(x); } template <> inline __host__ double log(double x) { return std::log(x); } template inline __host__ T pow(T x, T gamma) { return ck::type_convert( std::powf(ck::type_convert(x), ck::type_convert(gamma))); } template <> inline __host__ float pow(float x, float gamma) { return std::powf(x, gamma); } template <> inline __host__ double pow(double x, double gamma) { return std::pow(x, gamma); } template inline __host__ T expm1(T x) { return ck::type_convert(std::expm1f(ck::type_convert(x))); } template <> inline __host__ float expm1(float x) { return std::expm1f(x); } template <> inline __host__ double expm1(double x) { return std::expm1(x); } // math functions for the HIP kernel, some are implemented by calling hip builtin functions static inline __device__ float abs(float x) { return ::abs(x); }; static inline __device__ double abs(double x) { return ::abs(x); }; static inline __device__ int8_t abs(int8_t x) { int8_t sgn = x >> (8 - 1); return (x ^ sgn) - sgn; }; static inline __device__ int32_t abs(int32_t x) { int32_t sgn = x >> (32 - 1); return (x ^ sgn) - sgn; }; #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 static inline __device__ int4_t abs(int4_t x) { int4_t sgn = x >> (4 - 1); return (x ^ sgn) - sgn; }; #endif static inline __device__ half_t abs(half_t x) { uint16_t xx = ck::bit_cast(x); uint16_t abs_xx = xx & 0x7fff; half_t abs_x = ck::bit_cast(abs_xx); return abs_x; }; static inline __device__ bool isnan(float x) { return ::isnan(x); }; static inline __device__ bool isnan(double x) { return ::isnan(x); }; static inline __device__ bool isnan(int8_t x) { (void)x; return false; }; static inline __device__ bool isnan(int32_t x) { (void)x; return false; }; #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 static inline __device__ bool isnan(int4_t x) { (void)x; return false; }; #endif static inline __device__ bool isnan(half_t x) { uint16_t xx = ck::bit_cast(x); return (xx & 0x7FFF) > 0x7C00; }; static inline __device__ bool isnan(f8_t x) { return ck::fp8_is_nan(x); }; static inline __device__ half_t sqrt(half_t x) { return static_cast(__builtin_amdgcn_sqrtf(static_cast(x))); }; static inline __device__ float sqrt(float x) { return __builtin_amdgcn_sqrtf(x); }; static inline __device__ double sqrt(double x) { return __builtin_amdgcn_sqrt(x); }; template inline __device__ T tanh(T x) { return ck::type_convert(::tanhf(ck::type_convert(x))); }; template <> inline __device__ float tanh(float x) { return ::tanhf(x); }; template <> inline __device__ double tanh(double x) { return ::tanh(x); }; template inline __device__ T acos(T x) { return ck::type_convert(::acosf(ck::type_convert(x))); }; template <> inline __device__ float acos(float x) { return ::acosf(x); }; template <> inline __device__ double acos(double x) { return ::acos(x); }; template inline __device__ T neg(T x) { return ck::type_convert(-(ck::type_convert(x))); }; template <> inline __device__ float neg(float x) { return -x; }; template <> inline __device__ double neg(double x) { return -x; }; template <> inline __device__ int32_t neg(int32_t x) { return -x; }; template <> inline __device__ int8_t neg(int8_t x) { return -x; }; template <> inline __device__ half_t neg(half_t x) { return __hneg(static_cast<__half>(x)); }; template inline __device__ T atan(T x) { return ck::type_convert(::atanf(ck::type_convert(x))); }; template <> inline __device__ float atan(float x) { return ::atanf(x); }; template <> inline __device__ double atan(double x) { return ::atan(x); }; template inline __device__ T sin(T x) { return ck::type_convert(::sinf(ck::type_convert(x))); }; template <> inline __device__ float sin(float x) { return ::sinf(x); }; template <> inline __device__ double sin(double x) { return ::sin(x); }; template <> inline __device__ half_t sin(half_t x) { return hsin(static_cast<__half>(x)); }; template inline __device__ T asin(T x) { return ck::type_convert(::asinf(ck::type_convert(x))); }; template <> inline __device__ float asin(float x) { return ::asinf(x); }; template <> inline __device__ double asin(double x) { return ::asin(x); }; template inline __device__ T asinh(T x) { return ck::type_convert(::asinhf(ck::type_convert(x))); }; template <> inline __device__ float asinh(float x) { return ::asinhf(x); }; template <> inline __device__ double asinh(double x) { return ::asinh(x); }; template inline __device__ T acosh(T x) { return ck::type_convert(::acoshf(ck::type_convert(x))); }; template <> inline __device__ float acosh(float x) { return ::acoshf(x); }; template <> inline __device__ double acosh(double x) { return ::acosh(x); }; template inline __device__ T tan(T x) { return ck::type_convert(::tanf(ck::type_convert(x))); }; template <> inline __device__ float tan(float x) { return ::tanf(x); }; template <> inline __device__ double tan(double x) { return ::tan(x); }; template inline __device__ T atanh(T x) { return ck::type_convert(::atanhf(ck::type_convert(x))); }; template <> inline __device__ float atanh(float x) { return ::atanhf(x); }; template <> inline __device__ double atanh(double x) { return ::atanh(x); }; template inline __device__ T sinh(T x) { return ck::type_convert(::sinhf(ck::type_convert(x))); }; template <> inline __device__ float sinh(float x) { return ::sinhf(x); }; template <> inline __device__ double sinh(double x) { return ::sinh(x); }; template inline __device__ T ceil(T x) { return ck::type_convert(::ceilf(ck::type_convert(x))); }; template <> inline __device__ float ceil(float x) { return ::ceilf(x); }; template <> inline __device__ double ceil(double x) { return ::ceil(x); }; template <> inline __device__ half_t ceil(half_t x) { return hceil(static_cast<__half>(x)); }; template inline __device__ T cosh(T x) { return ck::type_convert(::coshf(ck::type_convert(x))); }; template <> inline __device__ float cosh(float x) { return ::coshf(x); }; template <> inline __device__ double cosh(double x) { return ::cosh(x); }; template inline __device__ T floor(T x) { return ck::type_convert(::floorf(ck::type_convert(x))); }; template <> inline __device__ float floor(float x) { return ::floorf(x); }; template <> inline __device__ double floor(double x) { return ::floor(x); }; template <> inline __device__ half_t floor(half_t x) { return hfloor(static_cast<__half>(x)); }; template inline __device__ T rcp(T x) { #if !CK_WORKAROUND_SWDEV_383542 return __frcp_rn(x); #else return __ocml_native_recip_f32(x); #endif }; template inline __device__ T exp(T x) { return ck::type_convert(__ocml_exp_f32(ck::type_convert(x))); }; template <> inline __device__ half_t exp(half_t x) { return hexp(static_cast<__half>(x)); }; template <> inline __device__ float exp(float x) { return __ocml_exp_f32(x); }; template <> inline __device__ double exp(double x) { return exp(x); }; template inline __device__ T log(T x) { return ck::type_convert(__logf(ck::type_convert(x))); }; template <> inline __device__ half_t log(half_t x) { return hlog(static_cast<__half>(x)); }; template <> inline __device__ float log(float x) { return __logf(x); }; template <> inline __device__ double log(double x) { return log(x); }; template inline __device__ T pow(T x, T gamma) { return ck::type_convert(powf(ck::type_convert(x), ck::type_convert(gamma))); }; template <> inline __device__ float pow(float x, float gamma) { return powf(x, gamma); }; template <> inline __device__ double pow(double x, double gamma) { return pow(x, gamma); }; template inline __device__ T expm1(T x) { return ck::type_convert(expm1f(ck::type_convert(x))); }; template <> inline __device__ float expm1(float x) { return expm1f(x); }; template <> inline __device__ double expm1(double x) { return expm1(x); }; } // namespace math } // namespace ck