diff --git a/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp index fadf4c194e..9cfe0f895c 100644 --- a/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_3d_grouped_bwd_xdlops.cpp @@ -86,14 +86,15 @@ struct CKArgs Do = ProblemInterpreter::GetOutputDepthDo(problem); Z = ProblemInterpreter::GetFilterDepthZ(problem); - // On a backward pass, out is in and in is out and this is silly - output = {G, N, C, Di, Hi, Wi}; - input = {G, N, K, Do, Ho, Wo}; + input = {G, N, C, Di, Hi, Wi}; + output = {G, N, K, Do, Ho, Wo}; weight = {G, K, C, Z, Y, X}; // miopen strides to CK strides - auto miopen_in_strides = problem.GetIn().GetStrides(); - auto miopen_out_strides = problem.GetOut().GetStrides(); + // On a backward pass, problem.GetIn() means y(or out), + // and problem.GetOut means x(or in) + auto miopen_in_strides = problem.GetOut().GetStrides(); + auto miopen_out_strides = problem.GetIn().GetStrides(); auto miopen_wei_strides = problem.GetWeights().GetStrides(); miopen_in_strides.insert(miopen_in_strides.begin(), C); miopen_out_strides.insert(miopen_out_strides.begin(), K); @@ -123,6 +124,29 @@ struct CKArgs template auto MakeArgPtr(const ConvPtr& conv_ptr, ConstData_t in, ConstData_t w, Data_t out) const { + +#if 0 // Leaving for debugging needs + std::cout << "y ptr = " << in << std::endl; + std::cout << "w ptr = " << w << std::endl; + std::cout << "x ptr = " << out << std::endl; + + auto print_vec = [](const char* name, const auto& vec) { + std::cout << name << " = [ "; + for(const auto& v : vec) + { + std::cout << v << ", "; + } + std::cout << "]\n"; + }; +#define PRINT_VEC(x) print_vec(#x, x); + + PRINT_VEC(output); + PRINT_VEC(out_strides); + PRINT_VEC(input); + PRINT_VEC(in_strides); + PRINT_VEC(weight); + PRINT_VEC(wei_strides); +#endif return conv_ptr->MakeArgumentPointer(in, w, {}, diff --git a/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp index deb8eb14d3..2ab65d9b1b 100644 --- a/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp @@ -84,14 +84,15 @@ struct CKArgs Do = ProblemInterpreter::GetOutputDepthDo(problem); Z = ProblemInterpreter::GetFilterDepthZ(problem); - // On a backward pass, out is in and in is out and this is silly - output = {G, N, C, Di, Hi, Wi}; - input = {G, N, K, Do, Ho, Wo}; + input = {G, N, C, Di, Hi, Wi}; + output = {G, N, K, Do, Ho, Wo}; weight = {G, K, C, Z, Y, X}; // miopen strides to CK strides - auto miopen_in_strides = problem.GetIn().GetStrides(); - auto miopen_out_strides = problem.GetOut().GetStrides(); + // On a backward pass, problem.GetIn() means y(or out), + // and problem.GetOut means x(or in) + auto miopen_in_strides = problem.GetOut().GetStrides(); + auto miopen_out_strides = problem.GetIn().GetStrides(); auto miopen_wei_strides = problem.GetWeights().GetStrides(); miopen_in_strides.insert(miopen_in_strides.begin(), C); miopen_out_strides.insert(miopen_out_strides.begin(), K); @@ -120,6 +121,28 @@ struct CKArgs template auto MakeArgPtr(const ConvPtr& conv_ptr, ConstData_t x, Data_t dw, ConstData_t dy) const { +#if 0 // Leaving for debugging needs + std::cout << "y ptr = " << dy << std::endl; + std::cout << "w ptr = " << dw << std::endl; + std::cout << "x ptr = " << x << std::endl; + + auto print_vec = [](const char* name, const auto& vec) { + std::cout << name << " = [ "; + for(const auto& v : vec) + { + std::cout << v << ", "; + } + std::cout << "]\n"; + }; +#define PRINT_VEC(x) print_vec(#x, x); + + PRINT_VEC(output); + PRINT_VEC(out_strides); + PRINT_VEC(input); + PRINT_VEC(in_strides); + PRINT_VEC(weight); + PRINT_VEC(wei_strides); +#endif return conv_ptr->MakeArgumentPointer(x, dw, dy, diff --git a/test/gtest/group_conv3d_bwd.cpp b/test/gtest/group_conv3d_bwd.cpp index 439dbb5eb7..10a63bb2a3 100644 --- a/test/gtest/group_conv3d_bwd.cpp +++ b/test/gtest/group_conv3d_bwd.cpp @@ -54,8 +54,9 @@ void SolverBwd(const miopen::TensorDescriptor& inputDesc, const auto tensors = miopen::ConvBwdTensors{outputDesc, output, wDesc, weight, inputDesc, input}; + // order for bwd data pass is dy, w, dx const auto problem = miopen::conv::ProblemDescription{ - inputDesc, wDesc, outputDesc, convDesc, miopen::conv::Direction::BackwardData}; + outputDesc, wDesc, inputDesc, convDesc, miopen::conv::Direction::BackwardData}; auto ctx = miopen::ExecutionContext{}; ctx.SetStream(&handle); diff --git a/test/gtest/group_conv3d_bwd.hpp b/test/gtest/group_conv3d_bwd.hpp index c59b48c500..9de7b4d18d 100644 --- a/test/gtest/group_conv3d_bwd.hpp +++ b/test/gtest/group_conv3d_bwd.hpp @@ -33,11 +33,11 @@ std::vector ConvTestConfigs() {1, 1, 1, 1, 4, 4, 1, 2, 2, 2, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, {1, 1, 1, 8, 8, 8, 1, 2, 2, 2, 0, 0, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}, {1, 1, 1, 8, 8, 8, 1, 2, 2, 2, 0, 0, 0, 2, 2, 2, 1, 1, 1, miopenConvolution}, - {1, 64, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {32, 128, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {1, 64, 32, 28, 28, 28, 16, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {16, 128, 16, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, {16, 128, 16, 28, 28, 28, 16, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {8, 128, 8, 28, 28, 28, 8, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {4, 128, 4, 28, 28, 28, 4, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {4, 128, 8, 28, 28, 28, 4, 2, 2, 2, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {4, 128, 4, 28, 28, 28, 8, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, {2, 128, 2, 28, 28, 28, 2, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}}; } diff --git a/test/gtest/group_conv3d_wrw.cpp b/test/gtest/group_conv3d_wrw.cpp index 8633b78b1e..7d3e31da5e 100644 --- a/test/gtest/group_conv3d_wrw.cpp +++ b/test/gtest/group_conv3d_wrw.cpp @@ -55,8 +55,9 @@ void SolverWrw(const miopen::TensorDescriptor& inputDesc, const auto tensors = miopen::ConvWrwTensors{outputDesc, output, inputDesc, input, wDesc, weight}; + // order for bwd weights pass is dy, dw, x const auto problem = miopen::conv::ProblemDescription{ - inputDesc, wDesc, outputDesc, convDesc, miopen::conv::Direction::BackwardWeights}; + outputDesc, wDesc, inputDesc, convDesc, miopen::conv::Direction::BackwardWeights}; auto ctx = miopen::ExecutionContext{}; ctx.SetStream(&handle); diff --git a/test/gtest/group_conv3d_wrw.hpp b/test/gtest/group_conv3d_wrw.hpp index 647004e45a..41d390d1f6 100644 --- a/test/gtest/group_conv3d_wrw.hpp +++ b/test/gtest/group_conv3d_wrw.hpp @@ -33,14 +33,14 @@ std::vector ConvTestConfigs() {1, 1, 1, 1, 4, 4, 1, 2, 2, 2, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, {1, 1, 1, 8, 8, 8, 1, 2, 2, 2, 0, 0, 0, 1, 1, 1, 1, 1, 1, miopenConvolution}, {1, 1, 1, 8, 8, 8, 1, 2, 2, 2, 0, 0, 0, 2, 2, 2, 1, 1, 1, miopenConvolution}, - {1, 64, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {2, 128, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {32, 128, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {8, 64, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {16, 64, 32, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {1, 64, 5, 28, 28, 28, 9, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {2, 128, 4, 28, 28, 28, 10, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {16, 128, 32, 28, 28, 28, 32, 2, 2, 2, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {8, 64, 32, 28, 28, 28, 16, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {16, 64, 16, 28, 28, 28, 32, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, {3, 48, 48, 28, 28, 28, 48, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {3, 48, 39, 28, 28, 28, 39, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, - {5, 120, 60, 28, 28, 28, 60, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}}; + {3, 48, 39, 28, 28, 28, 12, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}, + {5, 120, 10, 28, 28, 28, 30, 3, 3, 3, 1, 1, 1, 1, 1, 1, 1, 1, 1, miopenConvolution}}; } inline int SetTensorLayout(miopen::TensorDescriptor& desc)