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

Fix to the using of static_for in amd_buffer_addressing.hpp (#1337)

Browse Source 
* Add insert_dummy_dep_per_dword over-loading for length 64

* Fix insert_dummy_dep_per_dword and remove over-loading for length 64

* Remove blank lines

---------

Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: 37a347e380]
This commit is contained in:
Qianfeng
2024-06-13 16:12:20 +08:00
committed by GitHub
parent 2832eb1444
commit 9b0d87fe9a
1 changed files with 3 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
include/ck_tile/core/arch/amd_buffer_addressing.hpp
View File

@@ -552,8 +552,9 @@ namespace impl{
template<index_t N>
CK_TILE_DEVICE void insert_dummy_dep_per_dword(array<float, N>& b)
{
static_for<0, b.size(), 1>{}([&](auto i){
asm volatile(" " : : "v"(b.get(i)) : "memory");
constexpr auto kSize = remove_cvref_t<decltype(b)>::size();
static_for<0, kSize, 1>{}([&](auto i){
asm volatile(" " : : "v"(b.get(number<i>{})) : "memory");
});
}
#if 1