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

Add all device_grouped_conv_fwd_xdl_bf16_comp_instances

Browse Source
This commit is contained in:
John Shumway
2025-09-05 23:46:09 +00:00
parent cd1c1e0aff
commit 5da397b9ec
3 changed files with 149 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
experimental/builder/include/ck_tile/builder/conv_algorithm.hpp
View File

@@ -139,6 +139,7 @@ concept HasCBlockTransferInfo = requires(T t) {
enum class BlockGemmPipelineVersion
{
V1,
V3,
V4,
V5

1
experimental/builder/include/ck_tile/builder/conv_factory.hpp
View File

@@ -225,6 +225,7 @@ constexpr ck::BlockGemmPipelineVersion SetBlockGemmPipelineVersion()
{
switch(ALGORITHM.pipeline_version)
{
case BlockGemmPipelineVersion::V1: return ck::BlockGemmPipelineVersion::v1;
case BlockGemmPipelineVersion::V3: return ck::BlockGemmPipelineVersion::v3;
case BlockGemmPipelineVersion::V4: return ck::BlockGemmPipelineVersion::v4;
case BlockGemmPipelineVersion::V5: return ck::BlockGemmPipelineVersion::v5;

148
experimental/builder/test/test_conv_instances.cpp
View File

@@ -47,6 +47,7 @@ struct TestCase
// Test cases to drive the typed test suite.
constexpr std::array TEST_CASES = {
TestCase{
// double rate mfma instances on gfx950
.name = "ConvFwdXdlBf16CompInstances2x_0",
.algorithm =
{
@@ -65,6 +66,7 @@ constexpr std::array TEST_CASES = {
"2, 2, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>",
},
TestCase{
// Compute-friendly.
.name = "GroupedConvFwdXdlBf16CompInstance0",
.algorithm =
{
@@ -101,7 +103,7 @@ constexpr std::array TEST_CASES = {
"2, 2, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance1",
.name = "GroupedConvFwdXdlBf16CompInstance2",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 128, .n = 128, .k = 32}},
@@ -118,6 +120,150 @@ constexpr std::array TEST_CASES = {
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 128, 128, 32, Default, 32, 32, "
"2, 2, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance3",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 256, .k = 32}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V3,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, "
"4, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v3>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance4",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 256, .k = 32}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 4, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V5,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, "
"4, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v5>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance5",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 256, .n = 128, .k = 32}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V1,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 128, 32, Default, 32, 32, "
"2, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v1>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance7",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 128, .n = 256, .k = 32}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V1,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 128, 256, 32, Default, 32, 32, "
"2, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v1>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance8",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 128, .n = 128, .k = 64}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V1,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 128, 128, 64, Default, 32, 32, "
"2, 4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v1>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance9",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 128, .n = 64, .k = 64}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V3,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 128, 64, 64, Default, 32, 32, 2, "
"4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v3>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance9",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 64, .n = 128, .k = 64}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V3,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 64, 128, 64, Default, 32, 32, 2, "
"4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v3>",
},
TestCase{
.name = "GroupedConvFwdXdlBf16CompInstance9",
.algorithm =
{
.thread_block{.block_size = 256, .sub_matrix = {.m = 64, .n = 64, .k = 32}},
.tuning_params{.ak1 = 8, .bk1 = 8, .m_xdl_per_wave = 2, .n_xdl_per_wave = 4},
.block_transfer{
.thread_cluster_lengths_a = {.k0 = 4, .m = 64, .k1 = 1},
.thread_cluster_lengths_b = {.k0 = 4, .n = 64, .k1 = 1},
.thread_cluster_lengths_c =
{.m_block = 1, .m_wave_per_xdl = 32, .n_block = 1, .n_wave_per_xdl = 8},
},
.pipeline_version = ckb::BlockGemmPipelineVersion::V3,
},
.expected_type =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 64, 64, 32, Default, 32, 32, 2, "
"4, 8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v3>",
},
};
static constexpr int NUM_TEST_CASES = std::tuple_size_v<decltype(TEST_CASES)>;