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

Revert "Revert " Fp8 block scale quantization for fmha fwd (#3330)" (#3633)" (#3635)

Browse Source 
This reverts commit 723b7ce0be2884da131036301892bf9157f51876.

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 67f0b74ec6]
This commit is contained in:
ltqin
2026-01-24 01:03:22 +08:00
committed by GitHub
parent a0445aff5f
commit 90b3476006
14 changed files with 667 additions and 84 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
include/ck_tile/core/numeric/math.hpp
View File

@@ -37,6 +37,13 @@ struct scales
return lhs_ * rhs;
}
template <typename OtherScale>
CK_TILE_HOST_DEVICE constexpr auto operator*(OtherScale other) const
{
auto new_scale = lhs_ * other;
return scales<std::decay_t<decltype(new_scale)>>(new_scale);
}
private:
Scale lhs_;
};

12
include/ck_tile/core/utility/functional.hpp
View File

@@ -119,6 +119,18 @@ struct identity
}
};
// Similar to identity, but takes an additional index parameter as the first argument.
// The index is ignored and only the second argument (value) is forwarded.
// Useful for indexed element-wise operations where the functor signature requires an index.
struct idx_identity
{
template <typename I, typename T>
CK_TILE_HOST_DEVICE constexpr T&& operator()(I&& /*idx*/, T&& arg) const noexcept
{
return std::forward<T>(arg);
}
};
namespace detail {
// RemainLengths: sequence<...>