Merge pull request #3422 from ROCm/release-fixes/fix_3285

Produced solutions now contain original problem #3413

Author: vamovsik

src/include/miopen/problem.hpp

@@ -186,7 +186,8 @@ struct Problem
                                             const FindOptions& options,
                                             std::size_t max_solutions,
                                             const Buffers& buffers,
-                                            const ConvolutionDescriptor& conv_desc) const;
+                                            const ConvolutionDescriptor& conv_desc,
+                                            const Problem& original) const;
 
     std::vector<Solution> FindSolutionsImpl(Handle& handle,
                                             const FindOptions& options,

src/problem.cpp

@@ -179,7 +179,12 @@ Problem::FindSolutions(Handle& handle, const FindOptions& options, std::size_t m
     auto ret = std::visit(
         boost::hof::match(
             [&](const ConvolutionDescriptor& op_desc) {
-                return FindSolutionsImpl(handle, options, max_solutions, buffers, op_desc);
+                if(op_desc.mode == miopenTranspose)
+                    return MakeTransposed().FindSolutionsImpl(
+                        handle, options, max_solutions, buffers, op_desc, *this);
+                else
+                    return FindSolutionsImpl(
+                        handle, options, max_solutions, buffers, op_desc, *this);
             },
             [&](const SoftmaxDescriptor& op_desc) {
                 return FindSolutionsImpl(handle, options, max_solutions, buffers, op_desc);
@@ -460,7 +465,8 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
                                                  const FindOptions& options,
                                                  std::size_t max_solutions,
                                                  const Buffers& buffers,
-                                                 const ConvolutionDescriptor& conv_desc) const
+                                                 const ConvolutionDescriptor& conv_desc,
+                                                 const Problem& original) const
 {
     if(tensor_descriptors.size() != 3)
     {
@@ -477,21 +483,17 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
     const auto& w = buffers.at(miopenTensorConvolutionW);
     auto y        = buffers.at(miopenTensorConvolutionY);
 
-    const auto conv_problem =
-        conv_desc.mode == miopenTranspose ? MakeTransposed().AsConvolution() : AsConvolution();
-
-    std::size_t workspace_size;
-    Allocator::ManageDataPtr owned_workspace;
-    Data_t workspace;
-
     if(conv_desc.mode == miopenTranspose)
-    {
         std::swap(x, y);
-        std::swap(x_desc, y_desc);
-    }
+
+    const auto conv_problem = AsConvolution();
 
     ValidateGroupCount(x_desc, w_desc, conv_desc);
 
+    std::size_t workspace_size;
+    Allocator::ManageDataPtr owned_workspace;
+    Data_t workspace;
+
     if(options.preallocated_workspace)
     {
         workspace      = options.preallocated_workspace->buffer;
@@ -518,7 +520,7 @@ std::vector<Solution> Problem::FindSolutionsImpl(Handle& handle,
 
     for(auto& result : results)
     {
-        result.SetProblem({*this});
+        result.SetProblem({original});
 
         if(result.GetKernels().empty())
         {

test/CMakeLists.txt

@@ -930,6 +930,31 @@ add_custom_test(smoke_solver_ConvMlirIgemm_W GFX900_DISABLED GFX908_DISABLED GFX
     ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 ${IMPLICITGEMM_MLIR_ENV_W}
     COMMAND $<TARGET_FILE:test_conv2d> ${TEST_CONV_VERBOSE_W} --input 64 64  28 28 --weights 64  64  1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
 )
+
+add_custom_test(conv_transposed_fwd
+    COMMAND $<TARGET_FILE:test_conv2d> ${TEST_CONV_VERBOSE_F} --input 1 16 24 24 --weights 16 16 7 7 --pads_strides_dilations 3 3 1 1 1 1 --cmode transpose ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
+add_custom_test(conv_transposed_bwd
+    COMMAND $<TARGET_FILE:test_conv2d> ${TEST_CONV_VERBOSE_B} --input 64 64  28 28 --weights 64  64  1 1 --pads_strides_dilations 0 0 1 1 1 1 --cmode transpose ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
+add_custom_test(conv_transposed_wrw
+    COMMAND $<TARGET_FILE:test_conv2d> ${TEST_CONV_VERBOSE_W} --input 64 64  28 28 --weights 64  64  1 1 --pads_strides_dilations 0 0 1 1 1 1 --cmode transpose ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
+add_custom_test(conv_transposed_fwd_find_2
+    COMMAND $<TARGET_FILE:test_conv2d_find2> ${TEST_CONV_VERBOSE_F} --input 1 16 24 24 --weights 16 16 7 7 --pads_strides_dilations 3 3 1 1 1 1 --cmode transpose ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
+add_custom_test(conv_transposed_bwd_find_2
+    COMMAND $<TARGET_FILE:test_conv2d_find2> ${TEST_CONV_VERBOSE_B} --input 64 64  28 28 --weights 64  64  1 1 --pads_strides_dilations 0 0 1 1 1 1 --cmode transpose ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
+add_custom_test(conv_transposed_wrw_find_2
+    COMMAND $<TARGET_FILE:test_conv2d_find2> ${TEST_CONV_VERBOSE_W} --input 64 64  28 28 --weights 64  64  1 1 --pads_strides_dilations 0 0 1 1 1 1 --cmode transpose ${MIOPEN_TEST_FLAGS_ARGS}
+)
+
 #Dont install the tests if we are building each test separately
 #When the tests are built as a single unit, then we want to be
 #able to package them