ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-20 21:09:08 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit '4baa4c9fae0e56f1105d73a5d2484611d40886e0' into develop

Browse Source
This commit is contained in:
assistant-librarian[bot]
2025-12-03 20:13:41 +00:00
parent 834ee396bb
commit d0d2528beb
21 changed files with 2280 additions and 69 deletions
Download Patch File Download Diff File Expand all files Collapse all files

45
example/ck_tile/20_grouped_convolution/run_grouped_convolution_bwd_data_example.inc
View File

@@ -1,6 +1,9 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#pragma once
#include "ck_tile/ref/naive_grouped_conv_bwd_data_gpu.hpp"
template <ck_tile::index_t NDimSpatial,
typename ConvConfig,
typename Invoker,
@@ -185,7 +188,47 @@ int run_grouped_conv_bwd_data_example_with_layouts(
}
else if(arg_parser.get_int("v") == 2)
{
throw std::runtime_error("Unsupported gpu verification !!!");
// GPU reference verification
ck_tile::DeviceMem input_ref_dev_buf(input.get_element_space_size_in_bytes());
input_ref_dev_buf.SetZero();
// Launch GPU reference kernel
std::cout << "Run GPU reference kernel..." << std::endl;
ck_tile::naive_grouped_conv_bwd_data<NDimSpatial, InDataType, WeiDataType, OutDataType>(
reinterpret_cast<InDataType*>(input_ref_dev_buf.GetDeviceBuffer()),
reinterpret_cast<const WeiDataType*>(weight_dev_buf.GetDeviceBuffer()),
reinterpret_cast<const OutDataType*>(output_dev_buf.GetDeviceBuffer()),
conv_param.G_,
conv_param.N_,
conv_param.K_,
conv_param.C_,
conv_param.input_spatial_lengths_,
conv_param.filter_spatial_lengths_,
conv_param.output_spatial_lengths_,
conv_param.conv_filter_strides_,
conv_param.conv_filter_dilations_,
conv_param.input_left_pads_);
// Copy GPU reference result to host for comparison
ck_tile::HostTensor<InDataType> input_gpu_ref(in_g_n_c_wis_desc);
input_ref_dev_buf.FromDevice(input_gpu_ref.data());
const ck_tile::index_t GemmK = weight.get_element_size() / (conv_param.G_ * conv_param.K_);
const float max_accumulated_value =
*std::max_element(input_gpu_ref.mData.begin(), input_gpu_ref.mData.end());
const auto rtol_atol =
calculate_rtol_atol<InDataType, WeiDataType, AccDataType, OutDataType>(
GemmK, kbatch, max_accumulated_value);
pass = ck_tile::check_err(input,
input_gpu_ref,
"Error: Incorrect results!",
rtol_atol.at(ck_tile::number<0>{}),
rtol_atol.at(ck_tile::number<1>{}));
std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
<< " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
<< std::endl;
std::cout << "The GPU verification result is:" << (pass ? "correct" : "fail") << std::endl;
}
return pass;

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

@@ -1,6 +1,9 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#pragma once
#include "ck_tile/ref/naive_grouped_conv_bwd_weight_gpu.hpp"
template <ck_tile::index_t NDimSpatial,
typename ConvConfig,
typename Invoker,
@@ -185,7 +188,51 @@ int run_grouped_conv_bwd_weight_example_with_layouts(ck_tile::ArgParser& arg_par
}
else if(arg_parser.get_int("v") == 2)
{
throw std::runtime_error("Unsupported gpu verification !!!");
// GPU reference verification
ck_tile::DeviceMem weight_ref_dev_buf(weight.get_element_space_size_in_bytes());
weight_ref_dev_buf.SetZero();
// Launch GPU reference kernel
std::cout << "Run GPU reference kernel..." << std::endl;
ck_tile::naive_grouped_conv_bwd_weight<NDimSpatial, InDataType, WeiDataType, OutDataType>(
reinterpret_cast<const InDataType*>(input_dev_buf.GetDeviceBuffer()),
reinterpret_cast<WeiDataType*>(weight_ref_dev_buf.GetDeviceBuffer()),
reinterpret_cast<const OutDataType*>(output_dev_buf.GetDeviceBuffer()),
conv_param.G_,
conv_param.N_,
conv_param.K_,
conv_param.C_,
conv_param.input_spatial_lengths_,
conv_param.filter_spatial_lengths_,
conv_param.output_spatial_lengths_,
conv_param.conv_filter_strides_,
conv_param.conv_filter_dilations_,
conv_param.input_left_pads_);
// Copy GPU reference result to host for comparison
ck_tile::HostTensor<WeiDataType> weight_gpu_ref(wei_g_k_c_xs_desc);
weight_ref_dev_buf.FromDevice(weight_gpu_ref.data());
ck_tile::index_t GemmK = conv_param.N_;
for(ck_tile::index_t i = 0; i < NDimSpatial; ++i)
{
GemmK *= conv_param.output_spatial_lengths_[i];
}
const float max_accumulated_value =
*std::max_element(weight_gpu_ref.mData.begin(), weight_gpu_ref.mData.end());
const auto rtol_atol =
calculate_rtol_atol<InDataType, WeiDataType, AccDataType, OutDataType>(
GemmK, kbatch, max_accumulated_value);
pass = ck_tile::check_err(weight,
weight_gpu_ref,
"Error: Incorrect results!",
rtol_atol.at(ck_tile::number<0>{}),
rtol_atol.at(ck_tile::number<1>{}));
std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
<< " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
<< std::endl;
std::cout << "The GPU verification result is:" << (pass ? "correct" : "fail") << std::endl;
}
return pass;

6
example/ck_tile/20_grouped_convolution/run_grouped_convolution_fwd_bias_clamp_example.inc
View File

@@ -230,7 +230,11 @@ int run_grouped_conv_fwd_bias_clamp_example_with_layouts(
}
else if(arg_parser.get_int("v") == 2)
{
throw std::runtime_error("Unsupported gpu verification !!!");
// GPU verification for fused operation (Conv + Bias + Clamp) is complex
// For now, we only support GPU verification for basic convolution operations
// The bias+clamp fused variant can use CPU verification (-v=1) or no verification (-v=0)
throw std::runtime_error("GPU verification not yet supported for fused operations! Use "
"-v=1 for CPU verification.");
}
return pass;

46
example/ck_tile/20_grouped_convolution/run_grouped_convolution_fwd_example.inc
View File

@@ -3,6 +3,8 @@
#pragma once
#include "ck_tile/ref/naive_grouped_conv_fwd_gpu.hpp"
template <ck_tile::index_t NDimSpatial,
typename ConvConfig,
typename Invoker,
@@ -187,7 +189,49 @@ int run_grouped_conv_fwd_example_with_layouts(
}
else if(arg_parser.get_int("v") == 2)
{
throw std::runtime_error("Unsupported gpu verification !!!");
// GPU reference verification
ck_tile::DeviceMem output_ref_dev_buf(output.get_element_space_size_in_bytes());
output_ref_dev_buf.SetZero();
// GPU reference uses conv_param vectors directly (they are already long_index_t)
// Launch GPU reference kernel
std::cout << "Run GPU reference kernel..." << std::endl;
ck_tile::naive_grouped_conv_fwd<NDimSpatial, InDataType, WeiDataType, OutDataType>(
reinterpret_cast<const InDataType*>(input_dev_buf.GetDeviceBuffer()),
reinterpret_cast<const WeiDataType*>(weight_dev_buf.GetDeviceBuffer()),
reinterpret_cast<OutDataType*>(output_ref_dev_buf.GetDeviceBuffer()),
conv_param.G_,
conv_param.N_,
conv_param.K_,
conv_param.C_,
conv_param.input_spatial_lengths_,
conv_param.filter_spatial_lengths_,
conv_param.output_spatial_lengths_,
conv_param.conv_filter_strides_,
conv_param.conv_filter_dilations_,
conv_param.input_left_pads_);
// Copy GPU reference result to host for comparison
ck_tile::HostTensor<OutDataType> output_gpu_ref(out_g_n_k_wos_desc);
output_ref_dev_buf.FromDevice(output_gpu_ref.data());
const ck_tile::index_t GemmK = weight.get_element_size() / (conv_param.G_ * conv_param.K_);
const float max_accumulated_value =
*std::max_element(output_gpu_ref.mData.begin(), output_gpu_ref.mData.end());
const auto rtol_atol =
calculate_rtol_atol<InDataType, WeiDataType, AccDataType, OutDataType>(
GemmK, kbatch, max_accumulated_value);
pass = ck_tile::check_err(output,
output_gpu_ref,
"Error: Incorrect results!",
rtol_atol.at(ck_tile::number<0>{}),
rtol_atol.at(ck_tile::number<1>{}));
std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
<< " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
<< std::endl;
std::cout << "The GPU verification result is:" << (pass ? "correct" : "fail") << std::endl;
}
return pass;