ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-03 13:11:25 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix conv2d bwd data bug when filter is 1x1 and stride = 2 (#132)

Browse Source 
* fix bwd data filter1strid2 bug

* fichangeshort to ck::bhalf_t

* reset input to zero

Co-authored-by: ltqin <letaoqin@amd.com>
This commit is contained in:
ltqin
2022-03-21 23:53:23 +08:00
committed by GitHub
parent 9a17e7fbfd
commit b51808d7a5
5 changed files with 24 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
include/ck/tensor_operation/gpu/device/device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp
View File

@@ -459,6 +459,16 @@ struct DeviceConv2dBwdDataXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K
{
for(index_t i_xtilda = 0; i_xtilda < XTilda; ++i_xtilda)
{
// check slice is valid
const index_t Y = filter_spatial_lengths_[0];
const index_t X = filter_spatial_lengths_[1];
const auto YDotSlice = math::integer_divide_ceil(Y - i_ytilda, YTilda);
const auto XDotSlice = math::integer_divide_ceil(X - i_xtilda, XTilda);
if(YDotSlice * XDotSlice <= 0)
{
continue;
}
const auto descs = DeviceOp::MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(
N,
K,