From 813ad5a2ca3f67e188e54c9edf66dffc93283fd7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ville=20Pietil=C3=A4?= <> Date: Wed, 10 Dec 2025 10:34:38 -0500 Subject: [PATCH] Enable running int8 instances. --- .../tile_profiler_operation_registry.hpp | 6 ------ .../src/tile_profile_grouped_conv_fwd.cpp | 18 ++++++++++++------ .../gpu/tile_grouped_conv_fwd_factory.hpp | 2 +- 3 files changed, 13 insertions(+), 13 deletions(-) diff --git a/experimental/ck_tile_profiler/include/tile_profiler_operation_registry.hpp b/experimental/ck_tile_profiler/include/tile_profiler_operation_registry.hpp index 1f2acb93c4..d22960414c 100644 --- a/experimental/ck_tile_profiler/include/tile_profiler_operation_registry.hpp +++ b/experimental/ck_tile_profiler/include/tile_profiler_operation_registry.hpp @@ -56,12 +56,6 @@ class ProfilerOperationRegistry final { const auto found = entries_.find(name); - for (const auto& [key, value] : entries_) - { - // Debug output to trace available operations - std::cout << "Registered operation: " << key << " - " << value.description_ << std::endl; - } - if(found == end(entries_)) { return std::nullopt; diff --git a/experimental/ck_tile_profiler/src/tile_profile_grouped_conv_fwd.cpp b/experimental/ck_tile_profiler/src/tile_profile_grouped_conv_fwd.cpp index 283e6e7fc5..a78d0fd0a9 100644 --- a/experimental/ck_tile_profiler/src/tile_profile_grouped_conv_fwd.cpp +++ b/experimental/ck_tile_profiler/src/tile_profile_grouped_conv_fwd.cpp @@ -78,8 +78,8 @@ static void print_helper_msg() int tile_profile_grouped_conv_fwd(int argc, char* argv[]) { - // 8 for control, 1 for num_dim_spatial - if(argc < 9) + // 7 for control, 1 for num_dim_spatial + if(argc < 8) { print_helper_msg(); return 1; @@ -93,10 +93,12 @@ int tile_profile_grouped_conv_fwd(int argc, char* argv[]) const bool time_kernel = std::stoi(argv[7]); const int num_dim_spatial = std::stoi(argv[8]); - // 9 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial - if(argc != 8 + 1 + 4 + 6 * num_dim_spatial + 1) + // 7 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial + const int expected_num_args = 7 + 1 + 4 + 6 * num_dim_spatial + 1; + if(argc != expected_num_args) { - std::cout << argc << std::endl; + std::cout << "Received " << argc << " args"<< std::endl; + std::cout << "Expected " << expected_num_args << " args"<< std::endl; print_helper_msg(); return 1; } @@ -170,6 +172,10 @@ int tile_profile_grouped_conv_fwd(int argc, char* argv[]) { return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); } + if (data_type == ConvDataType::I8_I8_I8) + { + return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}); + } } if(num_dim_spatial == 3 && layout == ConvLayout::NHWGC_GKYXC_NHWGK) @@ -187,7 +193,7 @@ int tile_profile_grouped_conv_fwd(int argc, char* argv[]) return profile( I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{}); } - else if(data_type == ConvDataType::I8_I8_I8) + if(data_type == ConvDataType::I8_I8_I8) { return profile( I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}); diff --git a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_factory.hpp b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_factory.hpp index 49944f1a1e..5e0e91e2b6 100644 --- a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_factory.hpp +++ b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_factory.hpp @@ -65,7 +65,7 @@ void add_grouped_conv2d_fwd_bf16_instances_6(std::vector>& instances); // INT8 instances -void add_grouped_conv2d_fwd_i8_instances(std::vector>& instances); +void add_grouped_conv2d_fwd_int8_instances(std::vector>& instances); template