[NHWC] asm igemm xdlops fwd/bwd fp32/fp16 kernel

src/CMakeLists.txt

@@ -239,6 +239,9 @@ set( MIOpen_Source
     solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp
     solver/conv_asm_implicit_gemm_gtc_fwd.cpp
     solver/conv_asm_implicit_gemm_gtc_bwd.cpp
+    solver/conv_asm_implicit_gemm_gtc_perf_config.cpp
+    solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp
+    solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp
     solver/conv_direct_naive_conv_fwd.cpp
     solver/conv_direct_naive_conv_bwd.cpp
     solver/conv_direct_naive_conv_wrw.cpp

src/conv/invokers/impl_gemm_dynamic.cpp

@@ -416,5 +416,271 @@ MakeImplGemmDynamicBackwardDataInvokerFactory<solver::TunableImplicitGemmGTCDyna
     };
 }
 
+InvokerFactory MakeImplGemmDynamicForwardXdlopsNHWCInvokerFactory(
+    const ConvolutionContext& ctx,
+    const solver::PerformanceConfigAsmImplicitGemmGTCFwdXdlopsNHWC& config)
+{
+    const auto& conv_problem = ctx.conv_problem;
+    int hi                   = conv_problem.GetInHeight();
+    int wi                   = conv_problem.GetInWidth();
+    int n                    = conv_problem.GetInBatchSize();
+    int k                    = conv_problem.GetOutChannels();
+    int c                    = conv_problem.GetInChannels();
+    int ho                   = conv_problem.GetOutHeight();
+    int wo                   = conv_problem.GetOutWidth();
+    int stride_h             = conv_problem.GetKernelStrideH();
+    int stride_w             = conv_problem.GetKernelStrideW();
+    int dilation_h           = conv_problem.GetDilationH();
+    int dilation_w           = conv_problem.GetDilationW();
+    int pad_h                = conv_problem.GetPadH();
+    int pad_w                = conv_problem.GetPadW();
+    int y                    = conv_problem.GetWeightsHeight();
+    int x                    = conv_problem.GetWeightsWidth();
+    int group                = conv_problem.GetGroupCount();
+
+    uint32_t gemm_m = n * ho * wo;
+    uint32_t gemm_n = k / group;
+    magic_div_u32_t mdiv_0, mdiv_1, mdiv_2, mdiv_3, mdiv_4, mdiv_5;
+    uint32_t shift_pack_0, shift_pack_1;
+    uint32_t pack0 = 0;
+
+    mdiv_0 = magic_div_u32_gen((gemm_n + config.gemm_n_per_block - 1) / config.gemm_n_per_block);
+    mdiv_1 = magic_div_u32_gen(ho * wo);
+    mdiv_2 = magic_div_u32_gen(wo);
+    mdiv_3 = magic_div_u32_gen(((gemm_m + config.gemm_m_per_block - 1) / config.gemm_m_per_block) *
+                               ((gemm_n + config.gemm_n_per_block - 1) / config.gemm_n_per_block));
+
+    shift_pack_0 = magic_div_u32_pack_shift(mdiv_0.shift, mdiv_1.shift, mdiv_2.shift, mdiv_3.shift);
+    if(config.merge_e != 0)
+    {
+        mdiv_4       = magic_div_u32_gen(x * (c / group));
+        mdiv_5       = magic_div_u32_gen(c / group);
+        shift_pack_1 = magic_div_u32_pack_shift(mdiv_4.shift, mdiv_5.shift, 0, 0);
+
+        uint32_t s_move_slice_k_y = (config.gemm_k_per_block / (x * (c / group))) % y;
+        uint32_t s_move_slice_k_x = (config.gemm_k_per_block / (c / group)) % x;
+        uint32_t s_move_slice_k_c = config.gemm_k_per_block % (c / group);
+        y                         = static_cast<int>((s_move_slice_k_y << 24) | y);
+        x                         = static_cast<int>((s_move_slice_k_x << 24) | x);
+        c                         = static_cast<int>((s_move_slice_k_c << 24) | c);
+    }
+    else
+    {
+        mdiv_4       = magic_div_u32_gen(1);
+        mdiv_5       = magic_div_u32_gen(1);
+        shift_pack_1 = 0;
+    }
+
+    bool need_set_zero = config.gemm_k_global_split > 0;
+
+    std::vector<OpKernelArg> opArgs;
+    opArgs.emplace_back(0); // placeholder
+    opArgs.emplace_back(0); // placeholder
+    opArgs.emplace_back(0); // placeholder
+    opArgs.emplace_back(hi);
+    opArgs.emplace_back(wi);
+    opArgs.emplace_back(n);
+    opArgs.emplace_back(k / group);
+    opArgs.emplace_back(c / group);
+    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(group);
+    opArgs.emplace_back(mdiv_0.magic);
+    opArgs.emplace_back(mdiv_1.magic);
+    opArgs.emplace_back(mdiv_2.magic);
+    opArgs.emplace_back(mdiv_3.magic);
+    opArgs.emplace_back(mdiv_4.magic);
+    opArgs.emplace_back(mdiv_5.magic);
+    opArgs.emplace_back(shift_pack_0);
+    opArgs.emplace_back(shift_pack_1);
+    opArgs.emplace_back(config.gemm_k_global_split);
+    opArgs.emplace_back(pack0);
+
+    return [opArgs, need_set_zero](const std::vector<Kernel>& kernels) mutable {
+        return [=](const Handle& handle, const AnyInvokeParams& primitive_parameters) mutable {
+            decltype(auto) data_ctx = primitive_parameters.CastTo<conv::DataInvokeParams>();
+            const auto& tensors     = data_ctx.tensors;
+            const auto ker          = handle.Run(kernels[0]);
+            float elapsed           = 0;
+
+            opArgs[0] = OpKernelArg(tensors.in);
+            opArgs[1] = OpKernelArg(tensors.w);
+            opArgs[2] = OpKernelArg(tensors.out);
+
+            if(need_set_zero)
+            {
+                float zero = 0.f;
+                SetTensor(handle, tensors.outDesc, tensors.out, &zero);
+                if(handle.IsProfilingEnabled())
+                    elapsed += handle.GetKernelTime();
+            }
+
+            ker(opArgs);
+
+            if(handle.IsProfilingEnabled())
+            {
+                elapsed += handle.GetKernelTime();
+                handle.ResetKernelTime();
+                handle.AccumKernelTime(elapsed);
+            }
+        };
+    };
+}
+
+InvokerFactory MakeImplGemmDynamicBackwardDataXdlopsNHWCInvokerFactory(
+    const ConvolutionContext& ctx,
+    const solver::PerformanceConfigAsmImplicitGemmGTCBwdXdlopsNHWC& config)
+{
+    const auto& conv_problem = ctx.conv_problem;
+    int hi                   = conv_problem.GetOutHeight();
+    int wi                   = conv_problem.GetOutWidth();
+    int n                    = conv_problem.GetInBatchSize();
+    int k                    = conv_problem.GetInChannels();
+    int c                    = conv_problem.GetOutChannels();
+    int ho                   = conv_problem.GetInHeight();
+    int wo                   = conv_problem.GetInWidth();
+    int stride_h             = conv_problem.GetInHeight() > 1 ? conv_problem.GetKernelStrideH() : 1;
+    int stride_w             = conv_problem.GetInWidth() > 1 ? conv_problem.GetKernelStrideW() : 1;
+    int dilation_h = conv_problem.GetWeightsHeight() > 1 ? conv_problem.GetDilationH() : 1;
+    int dilation_w = conv_problem.GetWeightsWidth() > 1 ? conv_problem.GetDilationW() : 1;
+    int pad_h      = conv_problem.GetPadH();
+    int pad_w      = conv_problem.GetPadW();
+    int y          = conv_problem.GetWeightsHeight();
+    int x          = conv_problem.GetWeightsWidth();
+    int group      = conv_problem.GetGroupCount();
+
+    int gcd_stride_dilation_h = solver::gcd(stride_h, dilation_h);
+    int gcd_stride_dilation_w = solver::gcd(stride_w, dilation_w);
+    int y_tilda               = stride_h / gcd_stride_dilation_h;
+    int x_tilda               = stride_w / gcd_stride_dilation_w;
+
+    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;
+
+    uint32_t gemm_m = n * h_tilda_slice * w_tilda_slice;
+    uint32_t gemm_n = c / group;
+
+    magic_div_u32_t mdiv_x_tilda  = magic_div_u32_gen(x_tilda);
+    magic_div_u32_t mdiv_y_tilda  = magic_div_u32_gen(y_tilda);
+    magic_div_u32_t mdiv_group_mn = magic_div_u32_gen(
+        group * ((gemm_n + config.gemm_n_per_block - 1) / config.gemm_n_per_block) *
+        ((gemm_m + config.gemm_m_per_block - 1) / config.gemm_m_per_block));
+
+    magic_div_u32_t mdiv_0 =
+        magic_div_u32_gen((gemm_n + config.gemm_n_per_block - 1) / config.gemm_n_per_block);
+    magic_div_u32_t mdiv_1 =
+        magic_div_u32_gen(((gemm_n + config.gemm_n_per_block - 1) / config.gemm_n_per_block) *
+                          ((gemm_m + config.gemm_m_per_block - 1) / config.gemm_m_per_block));
+    magic_div_u32_t mdiv_2 = magic_div_u32_gen(config.nxe != 0 ? w_tilda_slice : wi);
+    magic_div_u32_t mdiv_3 = magic_div_u32_gen(h_tilda_slice * w_tilda_slice);
+    uint32_t shift_pack_0 =
+        magic_div_u32_pack_shift(mdiv_0.shift, mdiv_1.shift, mdiv_2.shift, mdiv_3.shift);
+
+    int dtile_iy = num_of_gemms > 1 ? static_cast<int>(mdiv_x_tilda.magic) : 0;
+    int dtile_ix = num_of_gemms > 1 ? static_cast<int>(mdiv_x_tilda.shift) : 0;
+    int dslice_y = num_of_gemms > 1 ? static_cast<int>(mdiv_y_tilda.magic) : y;
+    int dslice_x = num_of_gemms > 1 ? static_cast<int>(mdiv_y_tilda.shift) : x;
+    int dtile_h  = num_of_gemms > 1 ? static_cast<int>(mdiv_group_mn.magic) : h_tilda;
+    int dtile_w  = num_of_gemms > 1 ? static_cast<int>(mdiv_group_mn.shift) : w_tilda;
+
+    bool need_set_zero = false;
+    if(y < stride_h || x < stride_w || dilation_h != 1 || dilation_w != 1)
+        need_set_zero = true;
+    need_set_zero |= config.gemm_k_global_split > 0;
+
+    std::vector<OpKernelArg> opArgs;
+    opArgs.emplace_back(0); // placeholder
+    opArgs.emplace_back(0); // placeholder
+    opArgs.emplace_back(0); // placeholder
+    opArgs.emplace_back(hi);
+    opArgs.emplace_back(wi);
+    opArgs.emplace_back(n);
+    opArgs.emplace_back(k / group);
+    opArgs.emplace_back(c / group);
+    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(dilation_h / gcd_stride_dilation_h);
+    opArgs.emplace_back(dilation_w / gcd_stride_dilation_w);
+    opArgs.emplace_back(y_tilda);
+    opArgs.emplace_back(x_tilda);
+    opArgs.emplace_back(dtile_h);
+    opArgs.emplace_back(dtile_w);
+    opArgs.emplace_back(dslice_y);
+    opArgs.emplace_back(dslice_x);
+
+    opArgs.emplace_back(h_tilda_slice);
+    opArgs.emplace_back(w_tilda_slice);
+    opArgs.emplace_back(h_tilda_left);
+    opArgs.emplace_back(w_tilda_left);
+    opArgs.emplace_back(group);
+
+    opArgs.emplace_back(mdiv_0.magic);
+    opArgs.emplace_back(mdiv_1.magic);
+    opArgs.emplace_back(mdiv_2.magic);
+    opArgs.emplace_back(mdiv_3.magic);
+    opArgs.emplace_back(shift_pack_0);
+    opArgs.emplace_back(config.gemm_k_global_split);
+
+    return [opArgs, need_set_zero](const std::vector<Kernel>& kernels) mutable {
+        return [=](const Handle& handle, const AnyInvokeParams& primitive_parameters) mutable {
+            decltype(auto) data_ctx = primitive_parameters.CastTo<conv::DataInvokeParams>();
+            const auto& tensors     = data_ctx.tensors;
+            const auto ker          = handle.Run(kernels[0]);
+            float elapsed           = 0;
+
+            opArgs[0] = OpKernelArg(tensors.out);
+            opArgs[1] = OpKernelArg(tensors.w);
+            opArgs[2] = OpKernelArg(tensors.in);
+
+            if(need_set_zero)
+            {
+                float zero = 0.f;
+                SetTensor(handle, tensors.outDesc, tensors.out, &zero);
+                if(handle.IsProfilingEnabled())
+                    elapsed += handle.GetKernelTime();
+            }
+
+            ker(opArgs);
+
+            if(handle.IsProfilingEnabled())
+            {
+                elapsed += handle.GetKernelTime();
+                handle.ResetKernelTime();
+                handle.AccumKernelTime(elapsed);
+            }
+        };
+    };
+}
+
 } // namespace conv
 } // namespace miopen