ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-16 02:54:21 +00:00
Code Issues Packages Projects Releases Wiki Activity

Use compiler builtin directly in f_silu for float type

Browse Source
This commit is contained in:
Qianfeng Zhang
2025-04-25 15:04:41 +00:00
parent 4ae9acd712
commit 27f7ab4f2c
1 changed files with 8 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
example/ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp
View File

@@ -247,7 +247,14 @@ struct HstuAttentionFwdPipelineQRKSVS
const auto f_silu = [](CompDataType& x) {
const auto neg_one = ck_tile::type_convert<CompDataType>(-1.0f);
x = x * __builtin_amdgcn_rcpf(neg_one - exp(x));
if constexpr(std::is_same_v<CompDataType, float>)
{
x = x * __builtin_amdgcn_rcpf(neg_one - __expf(x));
}
else
{
x = x / (neg_one - exp(x));
}
};
using OaccBlockTileType = decltype(gemm_1.MakeCBlockTile());