ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 02:02:46 +00:00
Code Issues Packages Projects Releases Wiki Activity

[CK_TILE] Add CK Tile bwd weight profiler (#4797)

Browse Source 
## Motivation

To compare old CK and CK Tile, we need to extend the current CK profiler
to support running also CK Tile instance with the same API. In order to
have the same instance coverage in CK Tile compared to the old CK, I've
added code generation from old CK configurations to CK Tile instances
using the CK Builder.

## Technical Details

- The codegen python script for CK Tile fwd convs is extended to support
also bwd weight and bwd data.
- The generated instances are added to the CMake build (target
`device_grouped_conv_bwd_weight_tile_instance`s).
- A new profiler op (`grouped_conv_bwd_weight_tile`) has been added to
the CK Profiler.

---------

Co-authored-by: Ville Pietilä <>
Co-authored-by: Bartlomiej Kocot <barkocot@amd.com>
This commit is contained in:
Ville Pietilä
2026-03-04 23:49:42 +02:00
committed by GitHub
parent b6a914488f
commit 35bbebc3ce
68 changed files with 5194 additions and 196 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
example/ck_tile/20_grouped_convolution/conv_configs.hpp
View File

@@ -281,7 +281,9 @@ template <>
struct PipelineTypeTraits<ck_tile::GemmPipeline::BASIC_V1>
{
template <typename PipelineProblem>
using GemmPipeline = ck_tile::GemmPipelineAGmemBGmemCRegV1<PipelineProblem>;
using GemmPipeline =
ck_tile::GemmPipelineAGmemBGmemCRegV1<PipelineProblem,
ck_tile::GroupedConvUniversalPipelineAgBgCrPolicy>;
template <typename PipelineProblem>
using UniversalGemmPipeline = ck_tile::BaseGemmPipelineAGmemBGmemCRegV1<PipelineProblem>;
};
@@ -290,7 +292,9 @@ template <>
struct PipelineTypeTraits<ck_tile::GemmPipeline::BASIC_V2>
{
template <typename PipelineProblem>
using GemmPipeline = ck_tile::GemmPipelineAGmemBGmemCRegV2<PipelineProblem>;
using GemmPipeline =
ck_tile::GemmPipelineAGmemBGmemCRegV2<PipelineProblem,
ck_tile::GroupedConvUniversalPipelineAgBgCrPolicy>;
template <typename PipelineProblem>
using UniversalGemmPipeline = ck_tile::BaseGemmPipelineAGmemBGmemCRegV2<PipelineProblem>;
};
@@ -299,7 +303,9 @@ template <>
struct PipelineTypeTraits<ck_tile::GemmPipeline::MEMORY>
{
template <typename PipelineProblem>
using GemmPipeline = ck_tile::GemmPipelineAgBgCrMem<PipelineProblem>;
using GemmPipeline =
ck_tile::GemmPipelineAgBgCrMem<PipelineProblem,
ck_tile::GroupedConvUniversalPipelineAgBgCrPolicy>;
template <typename PipelineProblem>
using UniversalGemmPipeline = ck_tile::BaseGemmPipelineAgBgCrMem<PipelineProblem>;
};
@@ -308,7 +314,9 @@ template <>
struct PipelineTypeTraits<ck_tile::GemmPipeline::COMPUTE_V3>
{
template <typename PipelineProblem>
using GemmPipeline = ck_tile::GemmPipelineAgBgCrCompV3<PipelineProblem>;
using GemmPipeline =
ck_tile::GemmPipelineAgBgCrCompV3<PipelineProblem,
ck_tile::GroupedConvUniversalPipelineAgBgCrPolicy>;
template <typename PipelineProblem>
using UniversalGemmPipeline = ck_tile::BaseGemmPipelineAgBgCrCompV3<PipelineProblem>;
};