From db394d230ee2f8d3a9599c33aace59ae7cc3e4ef Mon Sep 17 00:00:00 2001 From: Po Yen Chen Date: Fri, 30 Aug 2024 10:40:00 +0800 Subject: [PATCH] Enable scratch memory workaround on ROCm 6.2 (#1486) Co-authored-by: carlushuang [ROCm/composable_kernel commit: 461ec98d78155ba95f01b62c4a575e4f1303a11d] --- include/ck_tile/core/config.hpp | 8 ++++++++ include/ck_tile/core/tensor/tile_window.hpp | 3 ++- 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index fa28aa2be9..ee47d136d0 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -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 diff --git a/include/ck_tile/core/tensor/tile_window.hpp b/include/ck_tile/core/tensor/tile_window.hpp index 04ed44201b..266d623c71 100644 --- a/include/ck_tile/core/tensor/tile_window.hpp +++ b/include/ck_tile/core/tensor/tile_window.hpp @@ -399,7 +399,8 @@ struct tile_window_with_static_distribution bottom_tensor_thread_coord, bool_constant{}, 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