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

Enable scratch memory workaround on ROCm 6.2 (#1486)

Browse Source 
Co-authored-by: carlushuang <carlus.huang@amd.com>

[ROCm/composable_kernel commit: 461ec98d78]
This commit is contained in:
Po Yen Chen
2024-08-30 10:40:00 +08:00
committed by GitHub
parent 7abc9cbe5f
commit db394d230e
2 changed files with 10 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
include/ck_tile/core/config.hpp
View File

@@ -156,6 +156,14 @@
#endif
#endif
#ifndef CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE
#if HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR == 2 && HIP_VERSION_PATCH >= 41133
#define CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE 1
#else
#define CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE 0
#endif
#endif
#ifndef CK_TILE_DEBUG_LOG
#define CK_TILE_DEBUG_LOG 0
#endif

3
include/ck_tile/core/tensor/tile_window.hpp
View File

@@ -399,7 +399,8 @@ struct tile_window_with_static_distribution
bottom_tensor_thread_coord,
bool_constant<oob_conditional_check>{},
pre_nop_);
#if CK_TILE_WORKAROUND_ROCM_6_1_SCRATCH_MEMORY_ISSUE
#if CK_TILE_WORKAROUND_ROCM_6_1_SCRATCH_MEMORY_ISSUE || \
CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE
asm volatile(
""); // this is starting from rocm-6.2, but same sympton, reuse this flag
#endif