mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-04-19 22:39:03 +00:00
37554c31e8
* [fix] align v3 gufusion pipeline * fix device kernel selection. * Add .co direct asm support by CK_USE_ASM_MOE_STAGE2_BLOCKSCALE * experimental optimization for scale load in blkscale gemm * Add asm for no-loop v3_128x128x128 * fix bugs * tune fp8 example * Update v1_128x128x128 to 2x2 instead of 4x1 * wip * add warmup to asm launch * wip2 * 16x16 function merged to moe * temp save, a performant version. * wip3 * Update .co binary to 16x16 * 16x16x128 correct; 64x64x128 failed * update * use mem_op::set when topk=1 * add mx fp8 b_preshuffle support, function not yet tested. * Spilt the fp4 target. Fix the known bugs. 128x128x128 sanity checked; remove prints * some fixes * fix update * remove some unnecessary hacky; enable 256x256x256 tilesize * update for function debug * Add pipeline v3. Have some runtime issue and register spill * Fix pipe v3 correctness issue * remove unnecessary hacky * clang format * fix a bug * fix the bug, functional test passed * tempsave; buggy at passed 4 e8m0 to scaled mfma * added fp4_bpreshuffle example, build failures * fixed some bugs * implement shuffled scale mxfp4gemm, blocker: opsel not effect * hotfix * fix bugs, build passed * (M, N, K)=(128, 128, 128) function failed. * temp save for gemm1. Function not ready * fix compile error. Gemm2 pass. Gemm1 WIP * fix bug for a lds read * update moe * Compile pass. Gemm1 function WIP * update moe * fix fp8; fix even/odd * tempsave * update moe * Revert "update" This reverts commit960b2bce1c. * Revert "use mem_op::set when topk=1" This reverts commitdef952a178. * Add v3 128x128x128_4x4_16x16.co for gfx950 * temp cmake flag suppression for aiter test * add code for mxfp4 gemm, blockscale not supported yet * gemm1 up-only pass. GU WIP * function pass with inline asm hacky * revert unexpected file change * updated and build passed * update CE elementOP * added code for debug * Gemm1 GUFusion function pass. Perf WIP * Fix fp8/bf8; remove duplicated code * disable the scheduler in v3; bring it back when compiler feature ready. * update moe v1 pipeline * Add gemm1 v1 32x128x128 * remove schedule barrier * updated * Fix fp8/bf8 B-row * mfma using asm, device result correct, host result need to check * gemm1 v3 64x128x128 debug * fix cpu ref * a/b thread_desc stride fix * Use random scale for init1 * 16x16x128 input size blockscale function passed * fix blockscale gemm bug * tempsave. Almost all instances passed. * v1 fix for mi350. * temp save * debug save * update debug * fix the bug, 128x128x256 tile function passed * v3 * rename moe block selector and pipeline * Add gemm1 v1 * Add gemm1 v1 to selector * added mx moe block v3 support, function passed * compile error fix * Improve the pipeline * Pack e8m0 as int32_t * v1 compile pass. Function not ready * debug synchronize issue over different GPU/ROCm * minor fix * Add profiler filter * Add f4 ckProfiler * Fix example compile error * Add f4 profiler examples * tempsave * v1 function pass. * v3 function pass * align file and function name * mx_moe_fp4 ready for aiter with clang-format. * modify the way we represent fp4 * generalize the pipeline scheduling. * init moe mx f4 scale shuffle * Cmakelist diable compiler-bound flags * mx_fp4 default parameter change * Moe blockscale gemm1&gemm2 asm support for aiter. Suppression cmkae flag til new compler. * update code * tempsave; modify the way we represent fp4 * generalize the pipeline scheduling. * Add gemm1 gfx942 .co support * updated code, build passed. * Update gemm2 asm with latest compiler flag * Fix mx f4 ckProfiler * Fix blockwise gemm mx v1 * lds conflict free + buffer load lds * Add gemm2 v3 64x128x128 * fix a, b scale loading bugs, a, b scale loading now correctly * Add gemm2 v3 64x128x128 * commit with debug info * fix fp4 profiler * Add mx fp4 pileline v1 instances * Fix v2 topk_weight cal. Add silu asm. * v2 tok_weight WIP * init mx fp4 B no preshuffle version * tempsave. compile pass, function wrong * enable fp4 moe no weigth preshuffle, function pass * update the TFlops calculation in the example * Add gemm2 64x128x128 asm. Fix BF16 ref. * fix 2 typos in fp4_preshuffle * Better kernel selection in device classes * correct preShuffleBuffer we should used packed k to do shuffle. * lds conflict free + buffer load lds * optimize offset math in dma * Fix fp4 ckProfiler * Fix MX MFMA tests * fix f4 pipeline issues * gemm1 func pass * update mx moe gemm1_bns tile size to 64x128x256 * update mx moe gemm1 gemm2 TF and BW calculation * fix typo * temp save * Fix example_gemm_mx build * rename the block pipeline * correct a typo in tail * Add rotating to mx examples * fix the correctness issue * Fix v1; use M padding * Add NT flag to B/BScale buffer * Merge gemm_mx_common.hpp * temp save, 4.4~4.5 * Fix 'Merge gemm_mx_common.hpp' * refactor the pipeline * Pad the M for scale buffer unconditionaly * update MX moe GEMM1 hotloopscheduling * change the gemm1 tile from 64x128x128 to 128x64x128 * Unconditional Ascale padding * Pad shuffled a scale only * pad ascale * add vmcnt guard for async copy * Profiler add f4 wp * Merge preshuffle device * Add more fp4 wp instances * Fix do_weight in gemm1. Fix cshuffle_datatype. Clang-format * Clang-format after 2 merges * Remove rocm6.3 workaround flags and macro * Fix fp8 config * Fix bf8 config * flag and barrier fix for copmiler branch MainOpSelV3 * Add fp8 profiler instances * Remove debug infos; Enable flags for blockscale f8 * No asm ver. for merging moe blocksale fp8 into mainline * update the flag name for f8blockscale * recover example * fix performance bug of bpreshuffle f8 gemm * clang format, remove single rate mfma restriction for f8 * remove single rate mfma restriction for f8 blockscale gemm * Fix moe blockscale gemm1 barrier 0x800 for new compiler * add pipeline v1 for MOE Gemm2 * Use v1 pipeline for example_moe_gemm2_xdl_mx_fp4_bns * Fix OOB; add MB96 instances * remove unnecessary files * fix the cmake issue * Enable splitk for mxfp4; clang format; * Generate random tensor values with multiple threads * Use packed_size_v for A/BPackedSize * Fix warning * Fix target_compile_options for disabled target on gfx942 * fix moe pki4 on gfx950 * doc the kGroup definition * Fix ThreadwiseTensorSliceTransfer_v4::Run (Fuse scale) * Refactor thread_copy_lds_direct_load; fix gfx942 direct lds load example; fix f16_pki4 example * Fix unknown compiler flag * fix two failed examples. * fix some failure tile size in gfx950 universal gemm. fix test_gemm_fp16 * workaround fix for test_gemm_f32; * We have very limited support for lds direct load if input matrix is not K major * fix test_gemm_splitk; * Fix compile for mx_mfma_op * add mfma selection logic for multipled_v3 * Clean up * Fix device gemm mx link error * improve the global atomic pattern * Revert unnecessary copyright updates * restore minimum_occupancy logic * Avoid data race in moe gemm2 ref * Build fp8 gemm_multiply_multiply and moe only on gfx94/95 * update the instance in device_mx_gemm * Resolve comments * Copyright 2025 * Remove unused code * fix library linking issue --------- Co-authored-by: OscarXu <huaiguxu@amd.com> Co-authored-by: lalala-sh <Jiaxing.Wen@amd.com> Co-authored-by: mtgu0705 <mtgu@amd.com> Co-authored-by: aska-0096 <haocwang@amd.com> Co-authored-by: Your Name <you@example.com> Co-authored-by: valarLip <340077269@qq.com> Co-authored-by: feifei14119 <feiw@amd.com> Co-authored-by: Lin, Qun <qlin@amd.com> Co-authored-by: Andriy Roshchenko <andriy.roshchenko@amd.com> Co-authored-by: joye <joye@amd.com> Co-authored-by: asleepzzz <hanwen.chang@amd.com>
246 lines
11 KiB
CMake
246 lines
11 KiB
CMake
include_directories(BEFORE
|
|
${PROJECT_SOURCE_DIR}/include
|
|
${PROJECT_SOURCE_DIR}/library/include
|
|
)
|
|
|
|
add_custom_target(examples)
|
|
|
|
|
|
# list of examples that are labelled as REGRESSION_EXAMPLE for make regression (runtime more than 30 seconds)
|
|
# all other tests are labelled as SMOKE_EXAMPLE
|
|
set(REGRESSION_EXAMPLES
|
|
example_sparse_embedding3_forward_layernorm
|
|
)
|
|
|
|
|
|
function(add_example_dependencies EXAMPLE_NAME FILE_NAME)
|
|
if(FILE_NAME)
|
|
add_dependencies(EXAMPLE_NAME FILE_NAME)
|
|
endif()
|
|
endfunction(add_example_dependencies EXAMPLE_NAME)
|
|
|
|
function(add_example_executable EXAMPLE_NAME FILE_NAME)
|
|
message(DEBUG "adding example ${EXAMPLE_NAME}")
|
|
set(result 1)
|
|
if(DEFINED DTYPES)
|
|
foreach(source IN LISTS FILE_NAME)
|
|
set(test 0)
|
|
if((source MATCHES "_fp16" OR source MATCHES "_f16") AND NOT "fp16" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_fp32" OR source MATCHES "_f32") AND NOT "fp32" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_fp64" OR source MATCHES "_f64") AND NOT "fp64" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_fp8" OR source MATCHES "_f8") AND NOT "fp8" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_bf8" OR source MATCHES "_bf8") AND NOT "bf8" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_bf16" OR source MATCHES "_b16") AND NOT "bf16" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_int8" OR source MATCHES "_i8") AND NOT "int8" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if(test EQUAL 1)
|
|
message(DEBUG "removing example source file ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
endif()
|
|
|
|
set(EX_TARGETS ${SUPPORTED_GPU_TARGETS})
|
|
|
|
#Do not build any DL examples if DL_KERNELS not set
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT DEFINED DL_KERNELS AND source MATCHES "_dl")
|
|
message(DEBUG "removing dl example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any DPP examples if DPP_KERNELS not set
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT DEFINED DPP_KERNELS AND source MATCHES "_dpp")
|
|
message(DEBUG "removing dpp example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any XDL examples if gfx9 targets are not on the list
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT EX_TARGETS MATCHES "gfx9" AND source MATCHES "_xdl")
|
|
message(DEBUG "removing xdl example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any WMMA examples if gfx11 targets are not on the list
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT EX_TARGETS MATCHES "gfx11" AND NOT EX_TARGETS MATCHES "gfx12" AND source MATCHES "_wmma")
|
|
message(DEBUG "removing wmma example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any microscaling examples if gfx950 target is not on the list
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT EX_TARGETS MATCHES "gfx950" AND source MATCHES "_mx")
|
|
message(DEBUG "removing microscaling example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any FP8 examples if CK_ENABLE_FP8 not set
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT DEFINED CK_ENABLE_FP8 AND source MATCHES "_fp8")
|
|
message(DEBUG "removing fp8 example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any BF8 examples if CK_ENABLE_BF8 not set
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT DEFINED CK_ENABLE_BF8 AND source MATCHES "_bf8")
|
|
message(DEBUG "removing bf8 example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
# Build fp8 gemm_multiply_multiply and moe only on gfx94/95
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT EX_TARGETS MATCHES "gfx94" AND NOT EX_TARGETS MATCHES "gfx95")
|
|
if (source MATCHES "fp8" AND source MATCHES "(gemm_multiply_multiply|moe)")
|
|
message(DEBUG "Skipping ${source} example for current target")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endif()
|
|
endforeach()
|
|
#only continue if there are some source files left on the list
|
|
if(FILE_NAME)
|
|
if(FILE_NAME MATCHES "_xdl" AND NOT FILE_NAME MATCHES "_pk_i4")
|
|
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
|
|
elseif(FILE_NAME MATCHES "_wmma")
|
|
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx950)
|
|
elseif(FILE_NAME MATCHES "_mx") #only build mx example for gfx950
|
|
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
|
|
elseif(FILE_NAME MATCHES "_pk_i4") #only build these examples for gfx942 and gfx950
|
|
message(DEBUG "trimming targets for ${FILE_NAME}")
|
|
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
|
|
endif()
|
|
set_source_files_properties(${FILE_NAME} PROPERTIES LANGUAGE HIP)
|
|
add_executable(${EXAMPLE_NAME} ${FILE_NAME})
|
|
target_link_libraries(${EXAMPLE_NAME} PRIVATE utility)
|
|
add_test(NAME ${EXAMPLE_NAME} COMMAND $<TARGET_FILE:${EXAMPLE_NAME}> ${ARGN})
|
|
set_property(TARGET ${EXAMPLE_NAME} PROPERTY HIP_ARCHITECTURES ${EX_TARGETS} )
|
|
add_dependencies(examples ${EXAMPLE_NAME})
|
|
add_dependencies(check ${EXAMPLE_NAME})
|
|
rocm_install(TARGETS ${EXAMPLE_NAME} COMPONENT examples)
|
|
set(result 0)
|
|
endif()
|
|
message(DEBUG "add_example returns ${result}")
|
|
if(result EQUAL 0 AND NOT "${EXAMPLE_NAME}" IN_LIST REGRESSION_EXAMPLES)
|
|
set_tests_properties(${EXAMPLE_NAME} PROPERTIES LABELS "SMOKE_TEST")
|
|
add_dependencies(smoke ${EXAMPLE_NAME})
|
|
elseif(result EQUAL 0 AND "${EXAMPLE_NAME}" IN_LIST REGRESSION_EXAMPLES)
|
|
set_tests_properties(${EXAMPLE_NAME} PROPERTIES LABELS "REGRESSION_TEST")
|
|
add_dependencies(regression ${EXAMPLE_NAME})
|
|
endif()
|
|
set(result ${result} PARENT_SCOPE)
|
|
endfunction(add_example_executable EXAMPLE_NAME)
|
|
|
|
function(add_example_dependencies EXAMPLE_NAME FILE_NAME)
|
|
if(result EQUAL 0)
|
|
add_dependencies(${EXAMPLE_NAME} ${FILE_NAME})
|
|
endif()
|
|
endfunction(add_example_dependencies EXAMPLE_NAME)
|
|
|
|
function(add_example_executable_no_testing EXAMPLE_NAME FILE_NAME)
|
|
message(DEBUG "adding example ${EXAMPLE_NAME}")
|
|
set(result 1)
|
|
if(DEFINED DTYPES)
|
|
foreach(source IN LISTS FILE_NAME)
|
|
set(test 0)
|
|
if((source MATCHES "_fp16" OR source MATCHES "_f16") AND NOT "fp16" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_fp32" OR source MATCHES "_f32") AND NOT "fp32" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_fp64" OR source MATCHES "_f64") AND NOT "fp64" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_fp8" OR source MATCHES "_f8") AND NOT "fp8" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_bf8" OR source MATCHES "_bf8") AND NOT "bf8" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_bf16" OR source MATCHES "_b16") AND NOT "bf16" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if((source MATCHES "_int8" OR source MATCHES "_i8") AND NOT "int8" IN_LIST DTYPES)
|
|
set(test 1)
|
|
endif()
|
|
if(test EQUAL 1)
|
|
message(DEBUG "removing example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
endif()
|
|
|
|
set(EX_TARGETS ${SUPPORTED_GPU_TARGETS})
|
|
|
|
#Do not build any DL examples if DL_KERNELS not set
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT DEFINED DL_KERNELS AND source MATCHES "_dl")
|
|
message(DEBUG "removing dl example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any XDL examples if gfx9 targets are not on the list
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT EX_TARGETS MATCHES "gfx9" AND source MATCHES "_xdl")
|
|
message(DEBUG "removing xdl example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#Do not build any WMMA examples if gfx11 targets are not on the list
|
|
foreach(source IN LISTS FILE_NAME)
|
|
if(NOT EX_TARGETS MATCHES "gfx11" AND NOT EX_TARGETS MATCHES "gfx12" AND source MATCHES "_wmma")
|
|
message(DEBUG "removing wmma example ${source} ")
|
|
list(REMOVE_ITEM FILE_NAME "${source}")
|
|
endif()
|
|
endforeach()
|
|
#only continue if there are some source files left on the list
|
|
if(FILE_NAME)
|
|
if(FILE_NAME MATCHES "_xdl")
|
|
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
|
|
elseif(FILE_NAME MATCHES "_wmma")
|
|
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx950)
|
|
endif()
|
|
set_source_files_properties(${FILE_NAME} PROPERTIES LANGUAGE HIP)
|
|
add_executable(${EXAMPLE_NAME} ${FILE_NAME})
|
|
target_link_libraries(${EXAMPLE_NAME} PRIVATE utility)
|
|
add_dependencies(examples ${EXAMPLE_NAME})
|
|
set_property(TARGET ${EXAMPLE_NAME} PROPERTY HIP_ARCHITECTURES ${EX_TARGETS} )
|
|
rocm_install(TARGETS ${EXAMPLE_NAME} COMPONENT examples)
|
|
set(result 0)
|
|
endif()
|
|
|
|
message(DEBUG "add_example returns ${result}")
|
|
set(result ${result} PARENT_SCOPE)
|
|
|
|
endfunction(add_example_executable_no_testing EXAMPLE_NAME)
|
|
|
|
function(example_compile_options EXAMPLE_NAME)
|
|
if(TARGET ${EXAMPLE_NAME})
|
|
target_compile_options(${EXAMPLE_NAME} ${ARGN})
|
|
endif()
|
|
endfunction(example_compile_options)
|
|
|
|
# add all example subdir
|
|
file(GLOB dir_list LIST_DIRECTORIES true *)
|
|
FOREACH(subdir ${dir_list})
|
|
if(IS_DIRECTORY "${subdir}" AND EXISTS "${subdir}/CMakeLists.txt")
|
|
add_subdirectory(${subdir})
|
|
ENDIF()
|
|
ENDFOREACH()
|