ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 18:42:06 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add transpose scale amax example (#1547)

Browse Source 
* Add transpose scale amax example

* fixes

* Tune reduce instance

[ROCm/composable_kernel commit: f21cda2536]
This commit is contained in:
Bartłomiej Kocot
2024-10-14 17:39:38 +02:00
committed by GitHub
parent 2117e76277
commit 87e4507543
4 changed files with 258 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp
View File

@@ -419,6 +419,12 @@ struct UnaryAbs
y = ck::math::abs(x);
};
template <>
__host__ __device__ void operator()(f8_t& y, const f8_t& x) const
{
y = ck::type_convert<f8_t>(ck::math::abs(ck::type_convert<float>(x)));
};
};
struct UnarySqrt

4
include/ck/utility/math_v2.hpp
View File

@@ -80,6 +80,8 @@ static inline __host__ bool isnan(half_t x)
return (xx & 0x7FFF) > 0x7C00;
};
static inline __host__ bool isnan(f8_t x) { return (x & 0x80); };
#ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
static inline __host__ bool isnan(int4_t x)
{
@@ -529,6 +531,8 @@ static inline __device__ bool isnan(half_t x)
return (xx & 0x7FFF) > 0x7C00;
};
static inline __device__ bool isnan(f8_t x) { return (x & 0x80); };
static inline __device__ half_t sqrt(half_t x)
{
return static_cast<half_t>(__builtin_amdgcn_sqrtf(static_cast<float>(x)));