ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 01:10:17 +00:00
Code Issues Packages Projects Releases Wiki Activity

temp fix buffer_store spill

Browse Source
This commit is contained in:
carlushuang
2024-03-15 22:56:41 +00:00
parent 04762d212b
commit ee397d0ab2
2 changed files with 4 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -413,7 +413,8 @@ struct buffer_store_if<8>
{
static_assert(sizeof(T) == 8);
auto save_exec = __builtin_amdgcn_read_exec();
using mbuf_t = fp32x2_t;
// TODO: ugly. rocm-6.0/6.1 seems neet bit_cast to same base type to avoid scratch
using mbuf_t = ext_vector_t<typename T::value_type::raw_type, T::size()>;
asm volatile("v_cmpx_le_u32 exec, 1, %5\n"
"buffer_store_dwordx2 %0, %1, %2, %3 offen offset:%4\n"
"s_mov_b64 exec %6"

4
include/ck_tile/core/numeric/half.hpp
View File

@@ -108,8 +108,8 @@ double fp16_to_double_hip(const fp16_hip_t& x) { return static_cast<double>(fp16
CK_TILE_HOST_DEVICE
fp16_hip_t float_to_fp16_hip(const float& x)
{
// return __float2half(x);
return static_cast<fp16_hip_t>(x);
return __float2half(x);
// return static_cast<fp16_hip_t>(x);
}
CK_TILE_HOST_DEVICE