diff --git a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.hpp b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.hpp new file mode 100644 index 0000000000..1e82d59d51 --- /dev/null +++ b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.hpp @@ -0,0 +1,49 @@ +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp" +#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_invoker.hpp" + +namespace ck_tile { +namespace ops { + +using BF16 = ck_tile::bfloat16_t; + +using DeviceOp2DBF16 = GroupedConvolutionBackwardWeightBaseInvoker<2, + NHWGC, + GKYXC, + NHWGK, + BF16, + BF16, + BF16, + PassThrough, + PassThrough, + PassThrough, + BF16, + BF16>; + +template +using tile_grouped_conv_bwd_weight_bf16_instances_opt = std::tuple< +// clang-format off + //#####################################| Num| InLayout| WeiLayout| OutLayout| InData| WeiData| OutData| In| Wei| Out| K-block| M-tile| N-tile | K-tile | M-warp| N-warp| K-warp| M-warp| N-warp| K-warp| Vector| Vector| Vector| Double| GEMM| + //#####################################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| smem| pipeline| + //#####################################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| buffer| version| + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + + // Skinny kernels. + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker + // clang-format on +>; + +} // namespace ops +} // namespace ck_tile diff --git a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_factory.hpp b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_factory.hpp index fa680930ba..0e776edd81 100644 --- a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_factory.hpp +++ b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_factory.hpp @@ -56,6 +56,7 @@ using DeviceOp2DF32 = GroupedConvolutionBackwardWeightBaseInvoker<2, // Forward declarations for instance factory functions void add_grouped_conv2d_bwd_weight_f16_instances(std::vector>& instances); void add_grouped_conv2d_bwd_weight_bf16_instances(std::vector>& instances); +void add_grouped_conv2d_bwd_weight_bf16_instances_opt(std::vector>& instances); template ) { add_grouped_conv2d_bwd_weight_bf16_instances(op_ptrs); + add_grouped_conv2d_bwd_weight_bf16_instances_opt(op_ptrs); } } } diff --git a/library/src/ck_tile/tensor_operation_instance/gpu/CMakeLists.txt b/library/src/ck_tile/tensor_operation_instance/gpu/CMakeLists.txt index 6331bec202..ac97cd0b97 100644 --- a/library/src/ck_tile/tensor_operation_instance/gpu/CMakeLists.txt +++ b/library/src/ck_tile/tensor_operation_instance/gpu/CMakeLists.txt @@ -11,6 +11,10 @@ add_library(ck_tile_grouped_conv_bwd_weight_bf16_instances STATIC tile_grouped_conv_bwd_weight_bf16_instances.cpp ) +add_library(ck_tile_grouped_conv_bwd_weight_bf16_instances_opt STATIC + tile_grouped_conv_bwd_weight_bf16_instances_opt.cpp +) + add_library(ck_tile_grouped_conv_fwd_fp16_instances STATIC tile_grouped_conv_fwd_fp16_instances.cpp ) @@ -29,6 +33,11 @@ target_include_directories(ck_tile_grouped_conv_bwd_weight_bf16_instances PRIVAT ${PROJECT_SOURCE_DIR}/library/include ) +target_include_directories(ck_tile_grouped_conv_bwd_weight_bf16_instances_opt PRIVATE + ${PROJECT_SOURCE_DIR}/include + ${PROJECT_SOURCE_DIR}/library/include +) + target_include_directories(ck_tile_grouped_conv_fwd_fp16_instances PRIVATE ${PROJECT_SOURCE_DIR}/include ${PROJECT_SOURCE_DIR}/library/include diff --git a/library/src/ck_tile/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.cpp b/library/src/ck_tile/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.cpp new file mode 100644 index 0000000000..8d315771ac --- /dev/null +++ b/library/src/ck_tile/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.cpp @@ -0,0 +1,20 @@ +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp" +#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.hpp" + +namespace ck_tile { +namespace ops { + +void add_grouped_conv2d_bwd_weight_bf16_instances_opt(std::vector>& instances) +{ + add_device_operation_instances(instances, + tile_grouped_conv_bwd_weight_bf16_instances_opt< + 2, + NHWGC, + GKYXC, + NHWGK>{}); +} + +} // namespace ops +} // namespace ck_tile diff --git a/profiler/ck_tile/src/CMakeLists.txt b/profiler/ck_tile/src/CMakeLists.txt index baa840cf59..26db40b9fa 100644 --- a/profiler/ck_tile/src/CMakeLists.txt +++ b/profiler/ck_tile/src/CMakeLists.txt @@ -41,6 +41,7 @@ set(DEVICE_INSTANCES "") if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx1[12]") list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_weight_fp16_instances) list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_weight_bf16_instances) + list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_weight_bf16_instances_opt) list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_fwd_fp16_instances) list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_fwd_bf16_instances) endif()