Respond to comments

Author: BalintCsala

projects/miopen/src/hip/utilocl.cpp

@@ -66,6 +66,8 @@ float Im2d2ColGPU(const Handle& handle,
 
     const int data_size_bound = c * in_h * in_w;
 
+    // NHWC variants use the common argument list below. NCHW launch parameters are
+    // shape-dependent, so rebuild that launch to supply its additional kernel arguments.
     if(!kernels.empty() && layoutNHWC)
     {
         auto kernel = kernels.front();
@@ -255,9 +257,8 @@ float Im2d2ColGPU(const Handle& handle,
             std::vector<size_t> vgd;
 
             const int bytes_per_pixel = c * get_data_size(type);
-            // TODO: Currently, we just pick whatever. In theory this
-            // should correspond with the cache line size for maximum
-            // efficiency.
+            // The many-channel kernel processes 16, 8, or 4 contiguous bytes per thread.
+            // Require the pixel stride to support the widest vectorization candidate.
             const int min_aligned_bytes = 16;
 
             bool selected = false;
@@ -297,14 +298,9 @@ float Im2d2ColGPU(const Handle& handle,
                     bytes_per_pixel % min_aligned_bytes == 0)
             {
 
-                // Arbitrary chosen: We expect the kernel to not use a lot of
-                // registers and have the maximum occupancy, therefore we can
-                // schedule at least 4 (with group_size_x hardcoded to 256) kernels
-                // on a single CU. Most AMD GPUs have in the order of 100-300 CUs,
-                // and we want to have at least one block for CU. This seems like
-                // a decent random value.
-                // TODO: This should be based on the GPU's CU count.
-                const int min_blocks = 256;
+                // Keep enough independent blocks to occupy every CU before reducing
+                // the amount of channel data processed by each thread.
+                const size_t min_blocks = handle.GetMaxComputeUnits();
 
                 int items_per_thread = 0;
                 int threads_per_ch   = 0;
@@ -399,7 +395,6 @@ float Im2d2ColGPU(const Handle& handle,
                 // OUTPUT PIXEL BASED VERSION
                 vgd = {
                     static_cast<size_t>(out_h) * out_w * group_size_x, 1, 1}; // outputpixel based
-                selected = true;
             }
 
             assert(vgd.size() == 3);

projects/miopen/src/solver/conv/gemm.cpp

@@ -712,9 +712,6 @@ bool GemmFwd1x1_0_1::IsApplicable(const ExecutionContext& context,
     if(!GemmFwdBase::IsApplicable(context, problem))
         return false;
 
-    if(!problem.IsLayoutDefault())
-        return false;
-
     decltype(auto) conv  = problem.GetConv();
     decltype(auto) wDesc = problem.GetWeights();
 
@@ -725,7 +722,7 @@ bool GemmFwd1x1_0_1::IsApplicable(const ExecutionContext& context,
     // for f8 on every architecture except gfx942. Grouped NHWC has no branch there at all.
     const auto nhwc_supported = problem.IsLayoutNHWC() && conv.group_count == 1 &&
                                 !problem.IsTensorsCasted() && !problem.IsFp8() && !problem.IsBfp8();
-    if(!problem.IsLayoutDefault() && !nhwc_supported)
+    if(!(problem.IsLayoutDefault() || nhwc_supported))
         return false;
 
     const auto spatial_dim = conv.GetSpatialDimension();
@@ -784,6 +781,7 @@ ConvSolution GemmFwd1x1_0_1::GetSolution(const ExecutionContext& context,
             tmp.strideA         = 0;
             tmp.strideB         = 0;
             tmp.strideC         = 0;
+            tmp.isColMajor      = false;
             tmp.m               = static_cast<int>(in_n * out_spatial_size);
             tmp.n               = static_cast<int>(wei_k);
             tmp.transA          = false;

projects/miopen/src/solver/conv/gemm_bwd.cpp

@@ -472,9 +472,6 @@ bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context,
     if(!GemmBwdBase::IsApplicable(context, problem))
         return false;
 
-    if(!problem.IsLayoutDefault())
-        return false;
-
     const auto& conv  = problem.GetConv();
     const auto& wDesc = problem.GetWeights();
 
@@ -487,7 +484,7 @@ bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context,
     // first while a_cast_type is filled from w.
     const auto nhwc_supported = problem.IsLayoutNHWC() && conv.group_count == 1 &&
                                 !problem.IsTensorsCasted() && !problem.IsFp8() && !problem.IsBfp8();
-    if(!problem.IsLayoutDefault() && !nhwc_supported)
+    if(!(problem.IsLayoutDefault() || nhwc_supported))
         return false;
 
     const auto spatial_dim = conv.GetSpatialDimension();
@@ -792,8 +789,7 @@ bool GemmBwdRest::IsApplicable(const ExecutionContext& context,
     if(!GemmBwdBase::IsApplicable(context, problem))
         return false;
 
-    // Everything below goes through Im2Col/Col2Im, which addresses dx channel-first.
-    if(!problem.IsLayoutDefault())
+    if(!(problem.IsLayoutDefault() || problem.IsLayoutNHWC()))
         return false;
 
     return !GemmBwd1x1_stride2{}.IsApplicable(context, problem) &&

projects/miopen/src/solver/conv/gemm_wrw.cpp

@@ -220,8 +220,6 @@ bool GemmWrw1x1_stride1::IsApplicable(const ExecutionContext& context,
 #if MIOPEN_USE_GEMM
     if(!GemmWrwBase::IsApplicable(context, problem))
         return false;
-    if(!problem.IsLayoutDefault())
-        return false;
 
     const auto& dwDesc = problem.GetWeights();
     const auto& conv   = problem.GetConv();
@@ -233,7 +231,7 @@ bool GemmWrw1x1_stride1::IsApplicable(const ExecutionContext& context,
     // for f8 on every architecture except gfx942. Grouped NHWC has no branch there at all.
     const auto nhwc_supported = problem.IsLayoutNHWC() && conv.group_count == 1 &&
                                 !problem.IsTensorsCasted() && !problem.IsFp8() && !problem.IsBfp8();
-    if(!problem.IsLayoutDefault() && !nhwc_supported)
+    if(!(problem.IsLayoutDefault() || nhwc_supported))
         return false;
 
     const auto wei_spatial =

projects/miopen/test/gtest/unit_conv_solver_GemmBwdRest.cpp

@@ -127,8 +127,6 @@ auto GetConvTestCases(miopenDataType_t datatype)
         TestCase{{1,  4,   7,   7},   {8,   2,   3,  3},  {1,  1},  {2,  1},  {1,  1},  2  ,datatype,miopenTensorNHWC},
         TestCase{{1,  4,   5,   5},   {4,   2,   3,  3},  {0,  0},  {1,  1},  {1,  1},  2  ,datatype,miopenTensorNHWC},
 
-        TestCase{{2, 16, 5, 5, 5}, {32, 16, 1, 1, 1}, {0, 0, 0}, {1, 1, 1}, {1, 1, 1}, datatype, datatype, datatype, miopenTensorNDHWC, miopenTensorNDHWC},
-        TestCase{{1, 64, 7, 7, 7}, {16, 64, 1, 1, 1}, {0, 0, 0}, {1, 1, 1}, {1, 1, 1}, datatype, datatype, datatype, miopenTensorNDHWC, miopenTensorNDHWC},
         TestCase{{1, 8, 14, 14, 14}, {16, 8, 3, 3, 3}, {0, 0, 0}, {1, 1, 1}, {1, 1, 1}, datatype, datatype, datatype, miopenTensorNDHWC, miopenTensorNDHWC},
         TestCase{{4, 3, 10, 10, 10}, {8, 3, 3, 3, 3}, {1, 1, 1}, {1, 1, 1}, {1, 1, 1}, datatype, datatype, datatype, miopenTensorNDHWC, miopenTensorNDHWC},
         TestCase{{2, 4, 16, 16, 16}, {8, 4, 3, 3, 3}, {1, 1, 1}, {2, 2, 2}, {1, 1, 1}, datatype, datatype, datatype, miopenTensorNDHWC, miopenTensorNDHWC},

projects/miopen/test/gtest/unit_conv_solver_GemmWrw1x1_stride1.cpp

@@ -112,7 +112,7 @@ TEST_P(CPU_UnitTestConvSolverGemmWrw1x1Stride1DevApplicabilityWrw_NONE, GemmWrw1
     this->RunTest(miopen::solver::conv::GemmWrw1x1_stride1{});
 };
 
-TEST(CPU_UnitTestConvSolverGemmWrw1x1Stride1Wrw_NONE, NHWCUsesUniversalSolver)
+TEST(CPU_UnitTestConvSolverGemmWrw1x1Stride1Wrw_NONE, NHWCUsesStride1Solver)
 {
     using TestCase = miopen::unit_tests::ConvTestCase;
 
@@ -123,8 +123,8 @@ TEST(CPU_UnitTestConvSolverGemmWrw1x1Stride1Wrw_NONE, NHWCUsesUniversalSolver)
     problem.SetupFloats(context);
     problem.SetupComputeType(context);
 
-    EXPECT_FALSE(miopen::solver::conv::GemmWrw1x1_stride1{}.IsApplicable(context, problem));
-    EXPECT_TRUE(miopen::solver::conv::GemmWrwUniversal{}.IsApplicable(context, problem));
+    EXPECT_TRUE(miopen::solver::conv::GemmWrw1x1_stride1{}.IsApplicable(context, problem));
+    EXPECT_FALSE(miopen::solver::conv::GemmWrwUniversal{}.IsApplicable(context, problem));
 }
 
 // Smoke tests