Commit 4bf3b87e authored by Po-Yen, Chen's avatar Po-Yen, Chen
Browse files

Use macro to toggle pipeline v2 optimization

parent 2a6d05f4
...@@ -100,7 +100,7 @@ struct GridwiseGemmPipeline_v2 ...@@ -100,7 +100,7 @@ struct GridwiseGemmPipeline_v2
// global read i + 2 // global read i + 2
b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf); b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf);
#if 1 #if defined(ENABLE_PIPELINE_V2_OPT)
__builtin_amdgcn_sched_group_barrier(0x020, 2, 0); // VMEM read __builtin_amdgcn_sched_group_barrier(0x020, 2, 0); // VMEM read
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
......
...@@ -20,7 +20,9 @@ add_library(gemm_standalone_xdl_fp16_instances STATIC ...@@ -20,7 +20,9 @@ add_library(gemm_standalone_xdl_fp16_instances STATIC
# instance/gemm_f16_tn_instance.cpp # instance/gemm_f16_tn_instance.cpp
# instance/gemm_f16_tt_instance.cpp # instance/gemm_f16_tt_instance.cpp
) )
set_source_files_properties(instance/gemm_f16_nn_instance.cpp PROPERTIES COMPILE_FLAGS "-save-temps -Wno-gnu-line-marker -amdgpu-enable-max-ilp-scheduling-strategy") set_source_files_properties(instance/gemm_f16_nn_instance.cpp PROPERTIES
COMPILE_FLAGS "-save-temps -Wno-gnu-line-marker -amdgpu-enable-max-ilp-scheduling-strategy"
COMPILE_DEFINITIONS ENABLE_PIPELINE_V2_OPT)
add_test_executable(test_gemm_standalone_xdl_fp16 gemm_standalone_xdl_fp16.cpp) add_test_executable(test_gemm_standalone_xdl_fp16 gemm_standalone_xdl_fp16.cpp)
target_link_libraries(test_gemm_standalone_xdl_fp16 PRIVATE gemm_standalone_xdl_fp16_instances utility) target_link_libraries(test_gemm_standalone_xdl_fp16 PRIVATE gemm_standalone_xdl_fp16_instances utility)
......
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