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

Disable the SLP vectorizer to prevent unnecessary wait (#1008)

Browse Source 
* Disable the SLP vectorizer to prevent unnecessary wait

* Add comment to the reason of adding flag

* Fix wording
This commit is contained in:
Po Yen Chen
2023-11-01 03:28:36 +08:00
committed by GitHub
parent 675b69782b
commit db4461c142
1 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
View File

@@ -108,13 +108,15 @@ if (ENABLE_PIPELINE_V2_OPT)
CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT=1
)
# TODO: The "-vectorize-slp=false" LLVM option is a workaround to prevent inefficient instruction scheduling
# caused by the SLP Vectorizer. Remove this option after fix the SLP Vectorizer issue.
# layout=NT
set_source_files_properties(device_gemm_xdl_f16_f16_f16/km_kn_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
COMPILE_OPTIONS ";;"
COMPILE_OPTIONS ";-mllvm;-vectorize-slp=false"
COMPILE_DEFINITIONS "${WAVES_PER_EU_DEFS};${IGLP_OPT_DEFS}")
# layout=NN
set_source_files_properties(device_gemm_xdl_f16_f16_f16/km_nk_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
COMPILE_OPTIONS ";;"
COMPILE_OPTIONS ";-mllvm;-vectorize-slp=false"
COMPILE_DEFINITIONS "${WAVES_PER_EU_DEFS};${IGLP_OPT_DEFS}")
# layout=TT
set_source_files_properties(device_gemm_xdl_f16_f16_f16/mk_kn_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES