ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-13 10:37:42 +00:00
Code Issues Packages Projects Releases Wiki Activity

Enable running int8 instances.

Browse Source
This commit is contained in:
Ville Pietilä
2025-12-10 10:34:38 -05:00
parent 224b2b5a94
commit 813ad5a2ca
3 changed files with 13 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
experimental/ck_tile_profiler/include/tile_profiler_operation_registry.hpp
View File

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

18
experimental/ck_tile_profiler/src/tile_profile_grouped_conv_fwd.cpp
View File

@@ -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{});

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

@@ -65,7 +65,7 @@ void add_grouped_conv2d_fwd_bf16_instances_6(std::vector<std::unique_ptr<DeviceO
void add_grouped_conv2d_fwd_f16_instances(std::vector<std::unique_ptr<DeviceOpFwd2DF16>>& instances);
// INT8 instances
void add_grouped_conv2d_fwd_i8_instances(std::vector<std::unique_ptr<DeviceOpFwd2DINT8>>& instances);
void add_grouped_conv2d_fwd_int8_instances(std::vector<std::unique_ptr<DeviceOpFwd2DINT8>>& instances);
template <ck_tile::index_t NumDimSpatial,
typename InLayout,