[igemm_dynamic] v4r1 bwd dynamic kernel

src/CMakeLists.txt

@@ -173,6 +173,7 @@ set( MIOpen_Source
     include/miopen/rnn_util.hpp
     include/miopen/bz2.hpp
     include/miopen/comgr.hpp
+    include/miopen/numeric.hpp
     md_graph.cpp
     mdg_expr.cpp
     conv/invokers/gcn_asm_1x1u.cpp
@@ -223,6 +224,7 @@ set( MIOpen_Source
     solver/conv_asm_implicit_gemm_v4r1_dynamic.cpp
     solver/conv_hip_implicit_gemm_bwd_data_v4r1_xdlops.cpp
     solver/conv_hip_implicit_gemm_v4r4_gen_xdlops_fwd_fp32.cpp
+    solver/conv_asm_implicit_gemm_bwd_v4r1_dynamic.cpp
     )
 
 list(APPEND MIOpen_Source tmp_dir.cpp binary_cache.cpp md5.cpp)

src/conv/invokers/impl_gemm_dynamic.cpp

@@ -3,23 +3,24 @@
 #include <miopen/algorithm.hpp>
 #include <miopen/handle.hpp>
 #include <miopen/tensor_ops.hpp>
-
+#include <miopen/numeric.hpp>
 #include <boost/any.hpp>
 
 namespace miopen {
 namespace conv {
 
-float CallImplicitGemmDynamic(const miopen::Handle& handle,
-                              const ConvolutionContext& ctx,
-                              ConstData_t src,
-                              Data_t dst,
-                              ConstData_t wei,
-                              const std::vector<KernelInvoke>& kernels)
+float CallImplGemmDynamicForward(const miopen::Handle& handle,
+                                 const ConvolutionContext& ctx,
+                                 ConstData_t src,
+                                 Data_t dst,
+                                 ConstData_t wei,
+                                 const std::vector<KernelInvoke>& kernels)
 {
     float elapsed = 0.0f;
 
     auto kernel = kernels[0];
     MIOPEN_LOG_I(kernel.GetName());
+
     bool kernel_is_1x1 = (kernel.GetName().find("igemm_v4r1_1x1_dynamic") == 0);
     // clang-format off
     int hi          = ctx.in_height;
@@ -56,6 +57,7 @@ float CallImplicitGemmDynamic(const miopen::Handle& handle,
     opArgs.emplace_back(dilation_w);
     opArgs.emplace_back(pad_h);
     opArgs.emplace_back(pad_w);
+    // clang-format off
     if(kernel_is_1x1)
     {
         opArgs.emplace_back(__pack0);
@@ -73,45 +75,198 @@ float CallImplicitGemmDynamic(const miopen::Handle& handle,
     return elapsed;
 }
 
-InvokerFactory MakeImplGemmDynamicDataInvokerFactory(const ConvolutionContext& ctx)
+
+float CallImplGemmDynamicBackwardData(const miopen::Handle& handle,
+                              const ConvolutionContext& ctx,
+                              ConstData_t src,
+                              Data_t dst,
+                              ConstData_t wei,
+                              const std::vector<KernelInvoke>& kernels)
 {
-    if(ctx.direction.IsForward())
+    float elapsed = 0.0f;
+
+    auto kernel = kernels[0];
+    MIOPEN_LOG_I(kernel.GetName());
+
+    // clang-format off
+    int hi          = ctx.out_height;
+    int wi          = ctx.out_width;
+    int n           = ctx.batch_sz;
+    int k           = ctx.n_inputs;
+    int c           = ctx.n_outputs;
+    int ho          = ctx.in_height;
+    int wo          = ctx.in_width;
+    int stride_h    = ctx.in_height > 1 ? ctx.kernel_stride_h : 1;
+    int stride_w    = ctx.in_width > 1 ? ctx.kernel_stride_w : 1;
+    int dilation_h  = ctx.kernel_size_h > 1? ctx.kernel_dilation_h : 1;
+    int dilation_w  = ctx.kernel_size_w > 1? ctx.kernel_dilation_w : 1;
+    int pad_h       = ctx.pad_h;
+    int pad_w       = ctx.pad_w;
+    int y           = ctx.kernel_size_h;
+    int x           = ctx.kernel_size_w;
+
+    int gcd_stride_dilation_h = gcd(stride_h, dilation_h);
+    int gcd_stride_dilation_w = gcd(stride_w, dilation_w);
+    int y_tilda     = stride_h / gcd_stride_dilation_h;
+    int x_tilda     = stride_w / gcd_stride_dilation_w;
+
+    int y_dot = (y +  y_tilda - 1) / y_tilda;
+    int x_dot = (x +  x_tilda - 1) / x_tilda;
+
+    int h_tilda     = ho + (dilation_h * (y - 1) + stride_h - 1) / stride_h;
+    int w_tilda     = wo + (dilation_w * (x - 1) + stride_w - 1) / stride_w;
+
+    int h_tilda_left = std::max(0, pad_h - dilation_h * (y_tilda - 1)) / stride_h;
+    int w_tilda_left = std::max(0, pad_w - dilation_w * (x_tilda - 1)) / stride_w;
+
+    int h_tilda_right = std::min(h_tilda, (pad_h + hi - 1 + stride_h - 1) / stride_h + 1);
+    int w_tilda_right = std::min(w_tilda, (pad_w + wi - 1 + stride_w - 1) / stride_w + 1);
+
+    int h_tilda_slice = h_tilda_right - h_tilda_left;
+    int w_tilda_slice = w_tilda_right - w_tilda_left;
+
+    int num_of_gemms  = x_tilda * y_tilda;
+
+    int dtile_iy    = 0;
+    int dtile_ix    = 0;
+    int dtile_dy    = dilation_h / gcd_stride_dilation_h;
+    int dtile_dx    = dilation_w / gcd_stride_dilation_w;
+    int dtile_y     = y_tilda;
+    int dtile_x     = x_tilda;
+    int dtile_h     = h_tilda;
+    int dtile_w     = w_tilda;
+    int dslice_y    = 0;
+    int dslice_x    = 0;
+    int dslice_h    = h_tilda_slice;
+    int dslice_w    = w_tilda_slice;
+    int dslice_h_left   = h_tilda_left;
+    int dslice_w_left   = w_tilda_left;
+    int __pack0     = 0;
+    // clang-format on
+
+    std::vector<OpKernelArg> opArgs;
+    opArgs.emplace_back(dst);
+    opArgs.emplace_back(wei);
+    opArgs.emplace_back(src);
+    opArgs.emplace_back(hi);
+    opArgs.emplace_back(wi);
+    opArgs.emplace_back(n);
+    opArgs.emplace_back(k);
+    opArgs.emplace_back(c);
+    opArgs.emplace_back(ho);
+    opArgs.emplace_back(wo);
+    opArgs.emplace_back(stride_h);
+    opArgs.emplace_back(stride_w);
+    opArgs.emplace_back(dilation_h);
+    opArgs.emplace_back(dilation_w);
+    opArgs.emplace_back(pad_h);
+    opArgs.emplace_back(pad_w);
+    opArgs.emplace_back(y);
+    opArgs.emplace_back(x);
+    opArgs.emplace_back(dtile_iy);
+    opArgs.emplace_back(dtile_ix);
+    opArgs.emplace_back(dtile_dy);
+    opArgs.emplace_back(dtile_dx);
+    opArgs.emplace_back(dtile_y);
+    opArgs.emplace_back(dtile_x);
+    opArgs.emplace_back(dtile_h);
+    opArgs.emplace_back(dtile_w);
+    opArgs.emplace_back(dslice_y);
+    opArgs.emplace_back(dslice_x);
+    opArgs.emplace_back(dslice_h);
+    opArgs.emplace_back(dslice_w);
+    opArgs.emplace_back(dslice_h_left);
+    opArgs.emplace_back(dslice_w_left);
+    opArgs.emplace_back(__pack0);
+
+    for(int gemm_id = 0; gemm_id < num_of_gemms; gemm_id++)
     {
-        return [ctx](const std::vector<Kernel>& kernels) {
-            return [=](const Handle& handle, const boost::any& primitive_parameters) {
-                const auto data_ctx = boost::any_cast<conv::DataInvokeParams>(primitive_parameters);
-                const auto& tensors = data_ctx.tensors;
-                auto kernel         = handle.Run(kernels[0]);
-                if(kernel.GetName().find("igemm_v4r1_dynamic") == 0 ||
-                   kernel.GetName().find("igemm_v4r1_1x1_dynamic") == 0)
+        int _dtile_iy          = gemm_id / x_tilda;
+        int _dtile_ix          = gemm_id % x_tilda;
+        int _y_dot_slice       = (_dtile_iy + 1) * y_dot <= y ? y_dot : y % y_dot;
+        int _x_dot_slice       = (_dtile_ix + 1) * x_dot <= x ? x_dot : x % x_dot;
+        int _gemm_k            = k * _y_dot_slice * _x_dot_slice;
+        bool is_gemm_not_empty = _gemm_k > 0;
+        opArgs[18]             = OpKernelArg(_dtile_iy);
+        opArgs[19]             = OpKernelArg(_dtile_ix);
+        opArgs[26]             = OpKernelArg(_y_dot_slice);
+        opArgs[27]             = OpKernelArg(_x_dot_slice);
+        if(is_gemm_not_empty)
+            kernel(opArgs);
+    }
+
+    if(handle.IsProfilingEnabled())
+        elapsed += handle.GetKernelTime();
+    return elapsed;
+}
+
+InvokerFactory MakeImplGemmDynamicForwardInvokerFactory(const ConvolutionContext& ctx)
+{
+    return [ctx](const std::vector<Kernel>& kernels) {
+        return [=](const Handle& handle, const boost::any& primitive_parameters) {
+            const auto data_ctx = boost::any_cast<conv::DataInvokeParams>(primitive_parameters);
+            const auto& tensors = data_ctx.tensors;
+            auto kernel         = handle.Run(kernels[0]);
+            if(kernel.GetName().find("igemm_v4r1_dynamic") == 0 ||
+               kernel.GetName().find("igemm_v4r1_1x1_dynamic") == 0)
+            {
+                std::vector<KernelInvoke> ks;
+                std::transform(kernels.begin(),
+                               kernels.end(),
+                               std::back_inserter(ks),
+                               [&](const Kernel& k) { return handle.Run(k); });
+                float elapsed = 0;
+                elapsed =
+                    CallImplGemmDynamicForward(handle, ctx, tensors.in, tensors.out, tensors.w, ks);
+                if(handle.IsProfilingEnabled())
                 {
-                    std::vector<KernelInvoke> ks;
-                    std::transform(kernels.begin(),
-                                   kernels.end(),
-                                   std::back_inserter(ks),
-                                   [&](const Kernel& k) { return handle.Run(k); });
-                    float elapsed = 0;
-                    elapsed       = CallImplicitGemmDynamic(
-                        handle, ctx, tensors.in, tensors.out, tensors.w, ks);
-                    if(handle.IsProfilingEnabled())
-                    {
-                        handle.ResetKernelTime();
-                        handle.AccumKernelTime(elapsed);
-                    }
+                    handle.ResetKernelTime();
+                    handle.AccumKernelTime(elapsed);
                 }
-                else
+            }
+            else
+            {
+                MIOPEN_THROW(
+                    "Error running dynamic implicit GEMM convolution (invalid kernel name " +
+                    kernel.GetName() + ")");
+            }
+        };
+    };
+}
+
+InvokerFactory MakeImplGemmDynamicBackwardDataInvokerFactory(const ConvolutionContext& ctx)
+{
+    return [ctx](const std::vector<Kernel>& kernels) {
+        return [=](const Handle& handle, const boost::any& primitive_parameters) {
+            const auto data_ctx = boost::any_cast<conv::DataInvokeParams>(primitive_parameters);
+            const auto& tensors = data_ctx.tensors;
+            auto kernel         = handle.Run(kernels[0]);
+            if(kernel.GetName().find("igemm_bwd_gtc") == 0)
+            {
+                std::vector<KernelInvoke> ks;
+                std::transform(kernels.begin(),
+                               kernels.end(),
+                               std::back_inserter(ks),
+                               [&](const Kernel& k) { return handle.Run(k); });
+                float elapsed = 0;
+
+                elapsed = CallImplGemmDynamicBackwardData(
+                    handle, ctx, tensors.in, tensors.out, tensors.w, ks);
+
+                if(handle.IsProfilingEnabled())
                 {
-                    MIOPEN_THROW(
-                        "Error running dynamic implicit GEMM convolution (invalid kernel name?)");
+                    handle.ResetKernelTime();
+                    handle.AccumKernelTime(elapsed);
                 }
-            };
+            }
+            else
+            {
+                MIOPEN_THROW(
+                    "Error running dynamic implicit GEMM convolution (invalid kernel name " +
+                    kernel.GetName() + ")");
+            }
         };
-    }
-    else
-    {
-        MIOPEN_THROW(
-            "Error running dynamic implicit GEMM convolution (currently only support forward)");
-    }
+    };
 }
 
 } // namespace conv

src/include/miopen/conv/invokers/impl_gemm_dynamic.hpp

@@ -35,17 +35,22 @@
 namespace miopen {
 namespace conv {
 
-// Beside used in invoker, currently this function is only called in RunAndMeasure() of dynamic
-// igemm solver
-// Remove this in the future when invoker is fully re-factored.
-float CallImplicitGemmDynamic(const miopen::Handle& handle,
-                              const ConvolutionContext& ctx,
-                              ConstData_t src,
-                              Data_t dst,
-                              ConstData_t wei,
-                              const std::vector<KernelInvoke>& kernels);
-
-InvokerFactory MakeImplGemmDynamicDataInvokerFactory(const ConvolutionContext& ctx);
+float CallImplGemmDynamicForward(const miopen::Handle& handle,
+                                 const ConvolutionContext& ctx,
+                                 ConstData_t src,
+                                 Data_t dst,
+                                 ConstData_t wei,
+                                 const std::vector<KernelInvoke>& kernels);
+
+float CallImplGemmDynamicBackwardData(const miopen::Handle& handle,
+                                      const ConvolutionContext& ctx,
+                                      ConstData_t src,
+                                      Data_t dst,
+                                      ConstData_t wei,
+                                      const std::vector<KernelInvoke>& kernels);
+
+InvokerFactory MakeImplGemmDynamicForwardInvokerFactory(const ConvolutionContext& ctx);
+InvokerFactory MakeImplGemmDynamicBackwardDataInvokerFactory(const ConvolutionContext& ctx);
 
 } // namespace conv
 } // namespace miopen

src/include/miopen/numeric.hpp

@@ -0,0 +1,64 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2020 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#ifndef GUARD_MLOPEN_NUMERIC_HPP
+#define GUARD_MLOPEN_NUMERIC_HPP
+
+#include <numeric>
+
+namespace miopen {
+
+template <typename T>
+T gcd(T x, T y)
+{
+    assert(!(x == 0 && y == 0));
+
+    if(x == y || x == 0)
+    {
+        return y;
+    }
+    else if(y == 0)
+    {
+        return x;
+    }
+    else if(x > y)
+    {
+        return gcd(x - y, y);
+    }
+    else
+    {
+        return gcd(x, y - x);
+    }
+}
+
+template <typename T, typename... Ys>
+T gcd(T x, Ys... ys)
+{
+    return gcd(x, gcd(ys...));
+}
+
+} // namespace miopen
+
+#endif

src/include/miopen/solver.hpp

@@ -1361,6 +1361,12 @@ struct ConvAsmImplicitGemmV4R1DynamicFwd_1x1 : SolverBase<ConvolutionContext>
                               float& elapsed_time) const;
 };
 
+struct ConvAsmImplicitGemmV4R1DynamicBwd : SolverBase<ConvolutionContext>
+{
+    bool IsApplicable(const ConvolutionContext&) const;
+    ConvSolution GetSolution(const ConvolutionContext&) const;
+};
+
 /// Holds common member functions for the Solvers which share the same
 /// "legacy exhaustive search" machinery.
 struct ConvOclDirectFwdLegacyExhaustiveSearch : SolverBase<ConvolutionContext>