mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-19 12:30:16 +00:00
8118d84f77
* Support 16x16 (MFMA, WMMA) and 32x32 (MFMA) tiles in fwd and bwd BlockDropout
Add comments with dropout implementation details
Fix performance regression of fwd+dropout
* Remove some usage of type punning (reinterpret_cast with ref or ptr) in Philox;
* "scalarize" seed and offset, they may come either from kernel args or from device memory
(presumably loaded with vector loads).
These changes help the compiler to procude more optimal code and reduce register spilling.
Use WarpGemmDispatcher instead of explicit WarpGemmMfma... to get CWarpDstrEncoding
Use code based on BlockDropout in BlockDropoutBwd
Refactor BlockDropout (fwd)
Implement BlockDropout (fwd) for WMMA
Originally BlockDropout only supported 32x32 tiles (IsWG32 = true),
this version supports 16x16 tiles.
If MPerBlock > MWarp * 16, it can generate numbers for two 16x16 tiles, similarly
to BlockDropoutBwd.
Implement BlockDropoutBwd for WMMA
Remove MakeRandValLds* functions unused in BlockDropoutBwd
Remove unused Run overload from BlockDropoutBwd
* Fix regression with philox seed and offset when they exceed 32-bit int
__builtin_amdgcn_readfirstlane works with 32-bit values, seed and offset
are 64-bit so they get truncated.
* Add F32 MFMA warp gemms
* Support f32 in fwd FMHA
* Implement transpose_vectors for 4-byte types (float)
* Fix unexpected implicit f32->uint32 cast in buffer_store<4>
__builtin_amdgcn_raw_buffer_store_b32 expects unsigned int but float was passed (implicitly casted to uint).
mbuf_t types in other buffer_store<> are changed for consistency.
* Support F32 in bwd FMHA
hdim = 256 is disabled for now because it uses too much memory on gfx90a
* Support Headdim = 48 (divisible by 16) in fwd
* Add fp32-specific receipts (800 and 801)
* Tune fwd tiles
* Tune bwd tiles
* Use small tiles only for small seqlen_q
* Fix after rebasing
* Fix selection of a fallback tile based on bm0
The assumption that the largest bm0 == 128 is not always true for
current fp32 tiles.
* Remove constraints and adjust filtering for fp32
Custom constraints are no longer needed because now the smallest tile
is selected automtically based on seqlen_q.
Filters related to qr_async_trload disabled valid fp32 tiles.
* Add fp32 tests
* Make splitkv and appendkv compile for fp32 only
There are no instances yet, but API still must compile when only fp32 is
requested.
* Remove unimportant f32 instances
* Add test_ck_tile_fmha_*_fp32 to REGRESSION_TESTS
* Replace magic numbers with a constant, improve comments for dropout
* Update changelog
* Fix condition that dq_acc must be set to zero when mask is used
The change was introduced in #2799
* Replace warp_uniform with recently added amd_wave_read_first_lane
* Add hdim = 96 and 192 to fwd
[ROCm/composable_kernel commit: 1edd250115]
296 lines
12 KiB
CMake
296 lines
12 KiB
CMake
include_directories(BEFORE
|
|
${PROJECT_SOURCE_DIR}/
|
|
${PROJECT_SOURCE_DIR}/profiler/include
|
|
)
|
|
|
|
include(gtest)
|
|
|
|
add_custom_target(tests)
|
|
|
|
# list of tests that are labelled as REGRESSION_TEST for make regression (runtime more than 30 seconds)
|
|
# all other tests are labelled as SMOKE_TEST
|
|
set(REGRESSION_TESTS
|
|
test_gemm_standalone_xdl_fp16
|
|
test_gemm_fp16
|
|
test_gemm_splitk
|
|
test_gemm_universal_wmma_fp16
|
|
test_gemm_universal_xdl_fp16
|
|
test_gemm_universal_streamk_fp16
|
|
test_gemm_universal_streamk_bf16
|
|
test_gemm_universal_streamk_fp8
|
|
test_batched_gemm_softmax_gemm_fp16
|
|
test_batched_gemm_softmax_gemm_permute_fp16
|
|
test_batched_gemm_bias_softmax_gemm_permute_fp16
|
|
test_batched_gemm_softmax_gemm_permute_bf16
|
|
test_batched_gemm_bias_softmax_gemm_permute_bf16
|
|
test_grouped_gemm_splitk
|
|
test_reduce_no_index
|
|
test_reduce_with_index
|
|
test_convnd_fwd
|
|
test_convnd_bwd_data
|
|
test_grouped_convnd_fwd
|
|
test_grouped_convnd_bwd_weight
|
|
test_softmax_rank3
|
|
test_softmax_rank4
|
|
test_batchnorm_fwd_rank_4
|
|
test_batchnorm_bwd_rank_4
|
|
test_grouped_convnd_bwd_data_xdl
|
|
test_conv_tensor_rearrange
|
|
test_gemm_mx
|
|
test_ck_tile_batched_transpose
|
|
test_ck_tile_fmha_bwd_fp32
|
|
test_ck_tile_fmha_bwd_bf16
|
|
test_ck_tile_fmha_bwd_fp16
|
|
test_ck_tile_fmha_fwd_fp32
|
|
test_ck_tile_fmha_fwd_bf16
|
|
test_ck_tile_fmha_fwd_fp16
|
|
test_ck_tile_fmha_fwd_fp8
|
|
)
|
|
|
|
function(add_test_executable TEST_NAME)
|
|
message(DEBUG "adding test ${TEST_NAME}")
|
|
set(result 1)
|
|
if(DEFINED DTYPES)
|
|
foreach(source IN LISTS ARGN)
|
|
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 test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
endif()
|
|
|
|
set(TEST_TARGETS ${SUPPORTED_GPU_TARGETS})
|
|
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT DEFINED DPP_KERNELS AND source MATCHES "_dpp")
|
|
message(DEBUG "removing dpp test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT DEFINED DL_KERNELS AND source MATCHES "_dl")
|
|
message(DEBUG "removing dl test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT TEST_TARGETS MATCHES "gfx9" AND NOT TEST_TARGETS MATCHES "gfx11" AND NOT TEST_TARGETS MATCHES "gfx12" AND source MATCHES "xdl")
|
|
message(DEBUG "removing xdl test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT TEST_TARGETS MATCHES "gfx11" AND NOT TEST_TARGETS MATCHES "gfx12" AND source MATCHES "wmma")
|
|
message(DEBUG "removing wmma test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
#only continue if there are some source files left on the list
|
|
if(ARGN)
|
|
if(ARGN MATCHES "_xdl")
|
|
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx10-3-generic)
|
|
elseif(ARGN MATCHES "_wmma")
|
|
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx950)
|
|
elseif(ARGN MATCHES "_smfmac")
|
|
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx908 gfx90a gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
|
|
endif()
|
|
set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
|
|
add_executable(${TEST_NAME} ${ARGN})
|
|
set_property(TARGET ${TEST_NAME} PROPERTY HIP_ARCHITECTURES ${TEST_TARGETS} )
|
|
target_link_libraries(${TEST_NAME} PRIVATE getopt::getopt)
|
|
add_test(NAME ${TEST_NAME} COMMAND $<TARGET_FILE:${TEST_NAME}>)
|
|
add_dependencies(tests ${TEST_NAME})
|
|
add_dependencies(check ${TEST_NAME})
|
|
rocm_install(TARGETS ${TEST_NAME} COMPONENT tests)
|
|
set(result 0)
|
|
endif()
|
|
message(DEBUG "add_test returns ${result}")
|
|
set(result ${result} PARENT_SCOPE)
|
|
if(result EQUAL 0 AND NOT "${TEST_NAME}" IN_LIST REGRESSION_TESTS)
|
|
set_tests_properties(${TEST_NAME} PROPERTIES LABELS "SMOKE_TEST")
|
|
add_dependencies(smoke ${TEST_NAME})
|
|
elseif(result EQUAL 0 AND "${TEST_NAME}" IN_LIST REGRESSION_TESTS)
|
|
set_tests_properties(${TEST_NAME} PROPERTIES LABELS "REGRESSION_TEST")
|
|
add_dependencies(regression ${TEST_NAME})
|
|
endif()
|
|
endfunction()
|
|
|
|
function(add_gtest_executable TEST_NAME)
|
|
message(DEBUG "adding gtest ${TEST_NAME}")
|
|
set(result 1)
|
|
if(DEFINED DTYPES)
|
|
foreach(source IN LISTS ARGN)
|
|
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 gtest ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
endif()
|
|
|
|
set(TEST_TARGETS ${SUPPORTED_GPU_TARGETS})
|
|
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT DEFINED DL_KERNELS AND source MATCHES "_dl")
|
|
message(DEBUG "removing dl test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT TEST_TARGETS MATCHES "gfx9" AND NOT TEST_TARGETS MATCHES "gfx1[12]" AND source MATCHES "xdl")
|
|
message(DEBUG "removing xdl test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT TEST_TARGETS MATCHES "gfx95" AND source MATCHES "mx_")
|
|
message(DEBUG "removing microscaling test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
|
|
foreach(source IN LISTS ARGN)
|
|
if(NOT TEST_TARGETS MATCHES "gfx11" AND NOT TEST_TARGETS MATCHES "gfx12" AND source MATCHES "wmma")
|
|
message(DEBUG "removing wmma test ${source} ")
|
|
list(REMOVE_ITEM ARGN "${source}")
|
|
endif()
|
|
endforeach()
|
|
|
|
#only continue if there are some source files left on the list
|
|
if(ARGN)
|
|
if(ARGN MATCHES "_xdl")
|
|
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx10-3-generic)
|
|
elseif(ARGN MATCHES "_wmma")
|
|
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx950)
|
|
elseif(ARGN MATCHES "_smfmac")
|
|
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx908 gfx90a gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
|
|
elseif(ARGN MATCHES "_mx") #only build mx example for gfx950
|
|
list(REMOVE_ITEM TEST_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)
|
|
endif()
|
|
set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
|
|
add_executable(${TEST_NAME} ${ARGN})
|
|
set_property(TARGET ${TEST_NAME} PROPERTY HIP_ARCHITECTURES ${TEST_TARGETS} )
|
|
add_dependencies(tests ${TEST_NAME})
|
|
add_dependencies(check ${TEST_NAME})
|
|
|
|
# suppress gtest warnings
|
|
target_compile_options(${TEST_NAME} PRIVATE -Wno-global-constructors -Wno-undef)
|
|
target_link_libraries(${TEST_NAME} PRIVATE gtest_main getopt::getopt)
|
|
add_test(NAME ${TEST_NAME} COMMAND $<TARGET_FILE:${TEST_NAME}>)
|
|
rocm_install(TARGETS ${TEST_NAME} COMPONENT tests)
|
|
set(result 0)
|
|
endif()
|
|
message(DEBUG "add_gtest returns ${result}")
|
|
set(result ${result} PARENT_SCOPE)
|
|
if(result EQUAL 0 AND NOT "${TEST_NAME}" IN_LIST REGRESSION_TESTS)
|
|
set_tests_properties(${TEST_NAME} PROPERTIES LABELS "SMOKE_TEST")
|
|
add_dependencies(smoke ${TEST_NAME})
|
|
elseif(result EQUAL 0 AND "${TEST_NAME}" IN_LIST REGRESSION_TESTS)
|
|
set_tests_properties(${TEST_NAME} PROPERTIES LABELS "REGRESSION_TEST")
|
|
add_dependencies(regression ${TEST_NAME})
|
|
endif()
|
|
endfunction()
|
|
|
|
add_compile_options(-Wno-c++20-extensions)
|
|
add_subdirectory(ck_tile)
|
|
add_subdirectory(magic_number_division)
|
|
add_subdirectory(space_filling_curve)
|
|
add_subdirectory(conv_util)
|
|
add_subdirectory(reference_conv_fwd)
|
|
add_subdirectory(gemm)
|
|
add_subdirectory(gemm_add)
|
|
add_subdirectory(gemm_layernorm)
|
|
add_subdirectory(gemm_multi_abd)
|
|
add_subdirectory(gemm_split_k)
|
|
add_subdirectory(gemm_universal)
|
|
add_subdirectory(gemm_b_scale)
|
|
add_subdirectory(gemm_universal_streamk)
|
|
add_subdirectory(gemm_reduce)
|
|
add_subdirectory(gemm_universal_reduce)
|
|
add_subdirectory(batched_gemm)
|
|
add_subdirectory(batched_gemm_reduce)
|
|
add_subdirectory(batched_gemm_gemm)
|
|
add_subdirectory(batched_gemm_softmax_gemm)
|
|
add_subdirectory(batched_gemm_softmax_gemm_permute)
|
|
add_subdirectory(grouped_gemm)
|
|
add_subdirectory(reduce)
|
|
add_subdirectory(convnd_fwd)
|
|
add_subdirectory(convnd_bwd_data)
|
|
add_subdirectory(grouped_convnd_fwd)
|
|
add_subdirectory(grouped_convnd_fwd_activation)
|
|
add_subdirectory(grouped_convnd_bwd_weight)
|
|
add_subdirectory(block_to_ctile_map)
|
|
add_subdirectory(softmax)
|
|
add_subdirectory(normalization_fwd)
|
|
add_subdirectory(normalization_bwd_data)
|
|
add_subdirectory(normalization_bwd_gamma_beta)
|
|
add_subdirectory(data_type)
|
|
add_subdirectory(elementwise_normalization)
|
|
add_subdirectory(batchnorm)
|
|
add_subdirectory(contraction)
|
|
add_subdirectory(pool)
|
|
add_subdirectory(batched_gemm_multi_d)
|
|
add_subdirectory(grouped_convnd_bwd_data)
|
|
add_subdirectory(conv_tensor_rearrange)
|
|
add_subdirectory(transpose)
|
|
add_subdirectory(permute_scale)
|
|
add_subdirectory(wrapper)
|
|
add_subdirectory(quantization)
|
|
if(SUPPORTED_GPU_TARGETS MATCHES "gfx11")
|
|
add_subdirectory(wmma_op)
|
|
endif()
|
|
if(SUPPORTED_GPU_TARGETS MATCHES "gfx942" OR SUPPORTED_GPU_TARGETS MATCHES "gfx950") # smfmac needs ROCm6.2
|
|
add_subdirectory(smfmac_op)
|
|
endif()
|
|
if(SUPPORTED_GPU_TARGETS MATCHES "gfx950")
|
|
add_subdirectory(mx_mfma_op)
|
|
add_subdirectory(gemm_mx)
|
|
endif()
|
|
add_subdirectory(position_embedding)
|
|
add_subdirectory(scatter_gather)
|