ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 01:27:34 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add new kernel instances.

Browse Source
This commit is contained in:
Ville Pietilä
2025-10-20 12:01:57 +00:00
parent 307ca52156
commit 2efd174b33
5 changed files with 81 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

49
library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.hpp Normal file
View File

@@ -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 <ck_tile::index_t NDimSpatial,
typename ALayout,
typename BLayout,
typename ELayout>
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<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 128, 128, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 128, 128, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
// Skinny kernels.
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
// clang-format on
>;
} // namespace ops
} // namespace ck_tile

2
library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_factory.hpp
View File

@@ -56,6 +56,7 @@ using DeviceOp2DF32 = GroupedConvolutionBackwardWeightBaseInvoker<2,
// Forward declarations for instance factory functions
void add_grouped_conv2d_bwd_weight_f16_instances(std::vector<std::unique_ptr<DeviceOp2DF16>>& instances);
void add_grouped_conv2d_bwd_weight_bf16_instances(std::vector<std::unique_ptr<DeviceOp2DBF16>>& instances);
void add_grouped_conv2d_bwd_weight_bf16_instances_opt(std::vector<std::unique_ptr<DeviceOp2DBF16>>& instances);
template <ck_tile::index_t NumDimSpatial,
typename InLayout,
@@ -117,6 +118,7 @@ struct DeviceOperationInstanceFactory<GroupedConvolutionBackwardWeightBaseInvoke
std::is_same_v<ComputeTypeB, ck_tile::bfloat16_t>)
{
add_grouped_conv2d_bwd_weight_bf16_instances(op_ptrs);
add_grouped_conv2d_bwd_weight_bf16_instances_opt(op_ptrs);
}
}
}

9
library/src/ck_tile/tensor_operation_instance/gpu/CMakeLists.txt
View File

@@ -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

20
library/src/ck_tile/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances_opt.cpp Normal file
View File

@@ -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<std::unique_ptr<DeviceOp2DBF16>>& instances)
{
add_device_operation_instances(instances,
tile_grouped_conv_bwd_weight_bf16_instances_opt<
2,
NHWGC,
GKYXC,
NHWGK>{});
}
} // namespace ops
} // namespace ck_tile

1
profiler/ck_tile/src/CMakeLists.txt
View File

@@ -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()