Skip to content

GitLab

Unverified Commit 461ec98d authored by Po Yen Chen's avatar Po Yen Chen Committed by GitHub
Browse files

Enable scratch memory workaround on ROCm 6.2 (#1486) 


Co-authored-by: default avatarcarlushuang <carlus.huang@amd.com>
parent 8107ee62
Show whitespace changes
Inline Side-by-side
Showing with 10 additions and 1 deletion
include/ck_tile/core/config.hpp
View file @ 461ec98d
...@@ -156,6 +156,14 @@ ...@@ -156,6 +156,14 @@
#endif #endif
#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 #ifndef CK_TILE_DEBUG_LOG
#define CK_TILE_DEBUG_LOG 0 #define CK_TILE_DEBUG_LOG 0
#endif #endif
......
include/ck_tile/core/tensor/tile_window.hpp
View file @ 461ec98d
...@@ -399,7 +399,8 @@ struct tile_window_with_static_distribution ...@@ -399,7 +399,8 @@ struct tile_window_with_static_distribution
bottom_tensor_thread_coord, bottom_tensor_thread_coord,
bool_constant<oob_conditional_check>{}, bool_constant<oob_conditional_check>{},
pre_nop_); 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( asm volatile(
""); // this is starting from rocm-6.2, but same sympton, reuse this flag ""); // this is starting from rocm-6.2, but same sympton, reuse this flag
#endif #endif
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment