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

[CK Tile] Fix building grouped conv examples in CK Tile (#2777)

Browse Source 
* Fix compilation of the grouped conv examples.

* Fix grouped conv bwd weight example output in CK Tile.
This commit is contained in:
Ville Pietilä
2025-09-05 09:14:21 +03:00
committed by GitHub
parent c6010f2953
commit 83f607e2a6
3 changed files with 5 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/ck_tile/20_grouped_convolution/run_grouped_convolution_bwd_weight_example.inc
View File

@@ -130,7 +130,7 @@ int run_grouped_conv_bwd_weight_example_with_layouts(
output_dev_buf.GetDeviceBuffer(), output_dev_buf.GetDeviceBuffer(),
kbatch); kbatch);
std::cout << "Run Grouped Conv Fwd kernel" << std::endl; std::cout << "Run Grouped Conv Bwd Weight kernel" << std::endl;
std::cout << "input: " << input.mDesc << std::endl; std::cout << "input: " << input.mDesc << std::endl;
std::cout << "weight: " << weight.mDesc << std::endl; std::cout << "weight: " << weight.mDesc << std::endl;
std::cout << "output: " << output.mDesc << std::endl; std::cout << "output: " << output.mDesc << std::endl;

5
include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp
View File

@@ -384,9 +384,8 @@ struct GroupedConvBwdDataKernelArgs
static constexpr index_t MaxGroupedGemmGroupsNum = 128; static constexpr index_t MaxGroupedGemmGroupsNum = 128;
using ABCGridDescs = using ABCGridDescs = remove_cvref_t<
remove_cvref_t<decltype(ConvToGemmTransformer{} decltype(ConvToGemmTransformer{}.MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(1))>;
.template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(1))>;
using AGridDescMK = remove_cvref_t<decltype(ABCGridDescs{}[number<0>{}])>; using AGridDescMK = remove_cvref_t<decltype(ABCGridDescs{}[number<0>{}])>;
using BGridDescNK = remove_cvref_t<decltype(ABCGridDescs{}[number<1>{}])>; using BGridDescNK = remove_cvref_t<decltype(ABCGridDescs{}[number<1>{}])>;

5
include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp
View File

@@ -254,9 +254,8 @@ struct GroupedConvBwdWeightKernelArgs
GemmBatch = args.G_; GemmBatch = args.G_;
} }
using ABCGridDescs = using ABCGridDescs = remove_cvref_t<
remove_cvref_t<decltype(ConvToGemmTransformer{} decltype(ConvToGemmTransformer{}.MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N())>;
.template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N())>;
using AGridDescMK = remove_cvref_t<decltype(ABCGridDescs{}[number<0>{}])>; using AGridDescMK = remove_cvref_t<decltype(ABCGridDescs{}[number<0>{}])>;
using BGridDescNK = remove_cvref_t<decltype(ABCGridDescs{}[number<1>{}])>; using BGridDescNK = remove_cvref_t<decltype(ABCGridDescs{}[number<1>{}])>;