From b1514f3b277b9b2ec72012c4d72eda101bc7aa5f Mon Sep 17 00:00:00 2001 From: Yi DING Date: Tue, 27 May 2025 09:51:20 +0800 Subject: [PATCH] Add operator/instance filters to ckProfiler (#2233) [ROCm/composable_kernel commit: 5727af98d1a70f9c4cc8aa7ac0cf6941690eeafb] --- profiler/src/CMakeLists.txt | 275 ++++++++++++++++++++---------------- 1 file changed, 156 insertions(+), 119 deletions(-) diff --git a/profiler/src/CMakeLists.txt b/profiler/src/CMakeLists.txt index 17c8c277eb..65dd704610 100644 --- a/profiler/src/CMakeLists.txt +++ b/profiler/src/CMakeLists.txt @@ -1,6 +1,16 @@ # ckProfiler -set(PROFILER_SOURCES - profiler.cpp +set(CK_PROFILER_OP_FILTER "" CACHE STRING "Filter for the operators to be profiled. Default is to include all") +set(CK_PROFILER_INSTANCE_FILTER "" CACHE STRING "Filter for the kernels instances to be profiled. Default is to be the same as the operator filter") +if (CK_PROFILER_OP_FILTER STREQUAL "") + set(CK_PROFILER_OP_FILTER ".+") +endif() +if (CK_PROFILER_INSTANCE_FILTER STREQUAL "") + set(CK_PROFILER_INSTANCE_FILTER ${CK_PROFILER_OP_FILTER}) +endif() +message(STATUS "CK_PROFILER_OP_FILTER: ${CK_PROFILER_OP_FILTER}") +message(STATUS "CK_PROFILER_INSTANCE_FILTER: ${CK_PROFILER_INSTANCE_FILTER}") + +set(PROFILER_OPS profile_gemm.cpp profile_reduce.cpp profile_groupnorm_bwd_data.cpp @@ -26,161 +36,188 @@ set(PROFILER_SOURCES if(SUPPORTED_GPU_TARGETS MATCHES "gfx9") if(DTYPES MATCHES "fp32" OR DTYPES MATCHES "fp64" OR NOT DEFINED DTYPES) - list(APPEND PROFILER_SOURCES profile_contraction_bilinear.cpp) - list(APPEND PROFILER_SOURCES profile_contraction_scale.cpp) + list(APPEND PROFILER_OPS profile_contraction_bilinear.cpp) + list(APPEND PROFILER_OPS profile_contraction_scale.cpp) endif() if(DTYPES MATCHES "fp16" OR NOT DEFINED DTYPES) - list(APPEND PROFILER_SOURCES profile_gemm_reduce.cpp) - list(APPEND PROFILER_SOURCES profile_batched_gemm_gemm.cpp) - list(APPEND PROFILER_SOURCES profile_batched_gemm_add_relu_gemm_add.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add_add_fastgelu.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add_fastgelu.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_gemm.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_streamk.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_fastgelu.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add_relu.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add_silu.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add_relu_add_layernorm.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_gemm_fixed_nk.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_gemm_fastgelu.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_gemm_tile_loop.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_gemm_multiply_tile_loop.cpp) + list(APPEND PROFILER_OPS profile_gemm_reduce.cpp) + list(APPEND PROFILER_OPS profile_batched_gemm_gemm.cpp) + list(APPEND PROFILER_OPS profile_batched_gemm_add_relu_gemm_add.cpp) + list(APPEND PROFILER_OPS profile_gemm_add.cpp) + list(APPEND PROFILER_OPS profile_gemm_add_add_fastgelu.cpp) + list(APPEND PROFILER_OPS profile_gemm_add_fastgelu.cpp) + list(APPEND PROFILER_OPS profile_grouped_gemm.cpp) + list(APPEND PROFILER_OPS profile_gemm_streamk.cpp) + list(APPEND PROFILER_OPS profile_gemm_fastgelu.cpp) + list(APPEND PROFILER_OPS profile_gemm_add_relu.cpp) + list(APPEND PROFILER_OPS profile_gemm_add_silu.cpp) + list(APPEND PROFILER_OPS profile_gemm_add_relu_add_layernorm.cpp) + list(APPEND PROFILER_OPS profile_grouped_gemm_fixed_nk.cpp) + list(APPEND PROFILER_OPS profile_grouped_gemm_fastgelu.cpp) + list(APPEND PROFILER_OPS profile_grouped_gemm_tile_loop.cpp) + list(APPEND PROFILER_OPS profile_grouped_gemm_multiply_tile_loop.cpp) endif() - list(APPEND PROFILER_SOURCES profile_gemm_multiply_add.cpp) - if(SUPPORTED_GPU_TARGETS MATCHES "gfx94") - list(APPEND PROFILER_SOURCES profile_gemm_multiply_multiply.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_multiply_multiply_wp.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_ab_scale.cpp) + list(APPEND PROFILER_OPS profile_gemm_multiply_add.cpp) + if(SUPPORTED_GPU_TARGETS MATCHES "gfx9[45]") + list(APPEND PROFILER_OPS profile_gemm_multiply_multiply.cpp) + list(APPEND PROFILER_OPS profile_gemm_multiply_multiply_wp.cpp) + list(APPEND PROFILER_OPS profile_gemm_ab_scale.cpp) endif() - list(APPEND PROFILER_SOURCES profile_batched_gemm.cpp) - list(APPEND PROFILER_SOURCES profile_batched_gemm_reduce.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_add_multiply.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_bias_add_reduce.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_splitk.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_b_scale.cpp) - list(APPEND PROFILER_SOURCES profile_batched_gemm_b_scale.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_universal_batched.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_universal_reduce.cpp) - list(APPEND PROFILER_SOURCES profile_gemm_universal_streamk.cpp) - list(APPEND PROFILER_SOURCES profile_conv_fwd_bias_relu.cpp) - list(APPEND PROFILER_SOURCES profile_conv_fwd_bias_relu_add.cpp) - list(APPEND PROFILER_SOURCES profile_conv_bwd_data.cpp) - list(APPEND PROFILER_SOURCES profile_conv_fwd.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_conv_fwd_outelementop.cpp) + list(APPEND PROFILER_OPS profile_batched_gemm.cpp) + list(APPEND PROFILER_OPS profile_batched_gemm_reduce.cpp) + list(APPEND PROFILER_OPS profile_gemm_add_multiply.cpp) + list(APPEND PROFILER_OPS profile_gemm_bias_add_reduce.cpp) + list(APPEND PROFILER_OPS profile_gemm_splitk.cpp) + list(APPEND PROFILER_OPS profile_gemm_b_scale.cpp) + list(APPEND PROFILER_OPS profile_batched_gemm_b_scale.cpp) + list(APPEND PROFILER_OPS profile_gemm_universal_batched.cpp) + list(APPEND PROFILER_OPS profile_gemm_universal_reduce.cpp) + list(APPEND PROFILER_OPS profile_gemm_universal_streamk.cpp) + list(APPEND PROFILER_OPS profile_conv_fwd_bias_relu.cpp) + list(APPEND PROFILER_OPS profile_conv_fwd_bias_relu_add.cpp) + list(APPEND PROFILER_OPS profile_conv_bwd_data.cpp) + list(APPEND PROFILER_OPS profile_conv_fwd.cpp) + list(APPEND PROFILER_OPS profile_grouped_conv_fwd_outelementop.cpp) endif() if(SUPPORTED_GPU_TARGETS MATCHES "gfx11" OR SUPPORTED_GPU_TARGETS MATCHES "gfx12" OR SUPPORTED_GPU_TARGETS MATCHES "gfx9") if(DTYPES MATCHES "fp16" OR NOT DEFINED DTYPES) - list(APPEND PROFILER_SOURCES profile_gemm_bilinear.cpp) + list(APPEND PROFILER_OPS profile_gemm_bilinear.cpp) endif() - list(APPEND PROFILER_SOURCES profile_gemm_universal.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_conv_fwd.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_conv_bwd_data.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_conv_bwd_weight.cpp) + list(APPEND PROFILER_OPS profile_gemm_universal.cpp) + list(APPEND PROFILER_OPS profile_grouped_conv_fwd.cpp) + list(APPEND PROFILER_OPS profile_grouped_conv_bwd_data.cpp) + list(APPEND PROFILER_OPS profile_grouped_conv_bwd_weight.cpp) endif() if(DL_KERNELS) - list(APPEND PROFILER_SOURCES profile_batched_gemm_multi_d.cpp) - list(APPEND PROFILER_SOURCES profile_grouped_conv_bwd_weight.cpp) + list(APPEND PROFILER_OPS profile_batched_gemm_multi_d.cpp) + list(APPEND PROFILER_OPS profile_grouped_conv_bwd_weight.cpp) endif() +set(PROFILER_SOURCES profiler.cpp) +foreach(SOURCE ${PROFILER_OPS}) + string(REGEX REPLACE "profile_(.+)\.cpp" "\\1" OP_NAME ${SOURCE}) + if (OP_NAME STREQUAL "") + message(FATAL_ERROR "Unexpected source file name: ${SOURCE}") + endif() + if("${OP_NAME}" MATCHES "${CK_PROFILER_OP_FILTER}") + list(APPEND PROFILER_SOURCES ${SOURCE}) + endif() +endforeach() +message(STATUS "ckProfiler sources: ${PROFILER_SOURCES}") + set(PROFILER_EXECUTABLE ckProfiler) add_executable(${PROFILER_EXECUTABLE} ${PROFILER_SOURCES}) target_compile_options(${PROFILER_EXECUTABLE} PRIVATE -Wno-global-constructors) # flags to compress the library if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 600241132) - message("Adding --offload-compress flag for ${PROFILER_EXECUTABLE}") + message(STATUS "Adding --offload-compress flag for ${PROFILER_EXECUTABLE}") target_compile_options(${PROFILER_EXECUTABLE} PRIVATE --offload-compress) endif() -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE utility getopt::getopt) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_normalization_fwd_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_normalization_bwd_data_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_normalization_bwd_gamma_beta_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_softmax_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_reduce_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batchnorm_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_pool2d_fwd_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_pool3d_fwd_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_avg_pool2d_bwd_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_avg_pool3d_bwd_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_max_pool_bwd_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_image_to_column_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_column_to_image_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_transpose_instance) -target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_permute_scale_instance) + +set(DEVICE_INSTANCES "") +list(APPEND DEVICE_INSTANCES device_gemm_instance) +list(APPEND DEVICE_INSTANCES device_normalization_fwd_instance) +list(APPEND DEVICE_INSTANCES device_normalization_bwd_data_instance) +list(APPEND DEVICE_INSTANCES device_normalization_bwd_gamma_beta_instance) +list(APPEND DEVICE_INSTANCES device_softmax_instance) +list(APPEND DEVICE_INSTANCES device_reduce_instance) +list(APPEND DEVICE_INSTANCES device_batchnorm_instance) +list(APPEND DEVICE_INSTANCES device_pool2d_fwd_instance) +list(APPEND DEVICE_INSTANCES device_pool3d_fwd_instance) +list(APPEND DEVICE_INSTANCES device_avg_pool2d_bwd_instance) +list(APPEND DEVICE_INSTANCES device_avg_pool3d_bwd_instance) +list(APPEND DEVICE_INSTANCES device_max_pool_bwd_instance) +list(APPEND DEVICE_INSTANCES device_image_to_column_instance) +list(APPEND DEVICE_INSTANCES device_column_to_image_instance) +list(APPEND DEVICE_INSTANCES device_transpose_instance) +list(APPEND DEVICE_INSTANCES device_permute_scale_instance) if(SUPPORTED_GPU_TARGETS MATCHES "gfx9") if(DTYPES MATCHES "fp32" OR DTYPES MATCHES "fp64" OR NOT DEFINED DTYPES) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_contraction_bilinear_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_contraction_scale_instance) + list(APPEND DEVICE_INSTANCES device_contraction_bilinear_instance) + list(APPEND DEVICE_INSTANCES device_contraction_scale_instance) endif() if(DTYPES MATCHES "fp16" OR NOT DEFINED DTYPES) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_add_fastgelu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_fastgelu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_gemm_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_add_relu_gemm_add_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_gemm_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_streamk_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_fastgelu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_relu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_silu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_relu_add_layernorm_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_gemm_fixed_nk_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_gemm_fastgelu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_gemm_tile_loop_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_add_fastgelu_instance) + list(APPEND DEVICE_INSTANCES device_gemm_fastgelu_instance) + list(APPEND DEVICE_INSTANCES device_batched_gemm_gemm_instance) + list(APPEND DEVICE_INSTANCES device_batched_gemm_add_relu_gemm_add_instance) + list(APPEND DEVICE_INSTANCES device_grouped_gemm_instance) + list(APPEND DEVICE_INSTANCES device_gemm_streamk_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_fastgelu_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_relu_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_silu_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_relu_add_layernorm_instance) + list(APPEND DEVICE_INSTANCES device_grouped_gemm_fixed_nk_instance) + list(APPEND DEVICE_INSTANCES device_grouped_gemm_fastgelu_instance) + list(APPEND DEVICE_INSTANCES device_grouped_gemm_tile_loop_instance) endif() - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_reduce_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_multiply_add_instance) - if(SUPPORTED_GPU_TARGETS MATCHES "gfx94") - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_multiply_multiply_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_multiply_multiply_wp_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_ab_scale_instance) + list(APPEND DEVICE_INSTANCES device_batched_gemm_instance) + list(APPEND DEVICE_INSTANCES device_batched_gemm_reduce_instance) + list(APPEND DEVICE_INSTANCES device_gemm_multiply_add_instance) + if(SUPPORTED_GPU_TARGETS MATCHES "gfx9[45]") + list(APPEND DEVICE_INSTANCES device_gemm_multiply_multiply_instance) + list(APPEND DEVICE_INSTANCES device_gemm_multiply_multiply_wp_instance) + list(APPEND DEVICE_INSTANCES device_gemm_ab_scale_instance) endif() - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_splitk_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_b_scale_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_b_scale_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_universal_batched_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_universal_reduce_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_universal_streamk_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_add_multiply_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_reduce_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_bias_add_reduce_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_bias_relu_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_fwd_bias_relu_add_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_fwd_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv1d_bwd_data_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv3d_bwd_data_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_conv2d_bwd_data_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_bwd_weight_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_weight_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_fwd_convscale_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_fwd_convinvscale_instance) + list(APPEND DEVICE_INSTANCES device_gemm_splitk_instance) + list(APPEND DEVICE_INSTANCES device_gemm_b_scale_instance) + list(APPEND DEVICE_INSTANCES device_batched_gemm_b_scale_instance) + list(APPEND DEVICE_INSTANCES device_gemm_universal_batched_instance) + list(APPEND DEVICE_INSTANCES device_gemm_universal_reduce_instance) + list(APPEND DEVICE_INSTANCES device_gemm_universal_streamk_instance) + list(APPEND DEVICE_INSTANCES device_gemm_add_multiply_instance) + list(APPEND DEVICE_INSTANCES device_gemm_reduce_instance) + list(APPEND DEVICE_INSTANCES device_gemm_bias_add_reduce_instance) + list(APPEND DEVICE_INSTANCES device_conv2d_fwd_instance) + list(APPEND DEVICE_INSTANCES device_conv2d_fwd_bias_relu_instance) + list(APPEND DEVICE_INSTANCES device_conv2d_fwd_bias_relu_add_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv1d_fwd_instance) + list(APPEND DEVICE_INSTANCES device_conv1d_bwd_data_instance) + list(APPEND DEVICE_INSTANCES device_conv3d_bwd_data_instance) + list(APPEND DEVICE_INSTANCES device_conv2d_bwd_data_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv1d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv2d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_fwd_convscale_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_fwd_convinvscale_instance) endif() if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx11" OR SUPPORTED_GPU_TARGETS MATCHES "gfx12") if(DTYPES MATCHES "fp16" OR NOT DEFINED DTYPES) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_bilinear_instance) + list(APPEND DEVICE_INSTANCES device_gemm_bilinear_instance) endif() - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_gemm_universal_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_fwd_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_data_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_bwd_data_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_fwd_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_gemm_universal_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_fwd_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv2d_bwd_data_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_bwd_data_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv2d_fwd_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_bwd_weight_instance) endif() if(DL_KERNELS) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_batched_gemm_multi_d_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv1d_bwd_weight_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_weight_instance) - target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv3d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_batched_gemm_multi_d_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv1d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv2d_bwd_weight_instance) + list(APPEND DEVICE_INSTANCES device_grouped_conv3d_bwd_weight_instance) endif() + +set(PROFILER_LIBS utility getopt::getopt) +foreach(LIB ${DEVICE_INSTANCES}) + string(REGEX REPLACE "device_(.+)_instance" "\\1" INSTANCE_NAME ${LIB}) + if (INSTANCE_NAME STREQUAL "") + message(FATAL_ERROR "Unexpected kernel instance name: ${LIB}") + endif() + if("${INSTANCE_NAME}" MATCHES "${CK_PROFILER_INSTANCE_FILTER}") + list(APPEND PROFILER_LIBS ${LIB}) + endif() +endforeach() +message(STATUS "ckProfiler libs: ${PROFILER_LIBS}") +target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE ${PROFILER_LIBS}) + rocm_install(TARGETS ${PROFILER_EXECUTABLE} COMPONENT profiler)