Merge branch 'develop' into simple_gemm_dl

include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp

@@ -110,10 +110,6 @@ __global__ void
     ignore = c_grid_desc_m0_m10_m11_n0_n10_n11;
     ignore = block_2_ctile_map;
     ignore = compute_ptr_offset_of_batch;
-
-    compute_ptr_offset_of_batch.GetAPtrOffset(0);
-    compute_ptr_offset_of_batch.GetBPtrOffset(0);
-    compute_ptr_offset_of_batch.GetCPtrOffset(0);
 #endif
 }
 

library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_dl_instance.hpp

@@ -54,7 +54,6 @@ using device_grouped_conv2d_fwd_dl_f16_instances = std::tuple<
            // ########################################|            |       |        |            |        |        |         |          |          |          |   Operation|    Operation|     Operation| Specialization|                  |      |      |      |      |   |           |           |       |              |              |        K0_M0_M1_K1|          K0_M0_M1_K1|   ArrangeOrder|          Order| Lengths_K0_M0_M1_K1| ContiguousDimOrder| Lengths_K0_M0_M1_K1|        K0_N0_N1_K1|          K0_N0_N1_K1|   ArrangeOrder|          Order| Lengths_K0_N0_N1_K1| ContiguousDimOrder| Lengths_K0_N0_N1_K1|               Order|                |                   |
            // ########################################|            |       |        |            |        |        |         |          |          |          |            |             |              |               |                  |      |      |      |      |   |           |           |       |              |              |                   |                     |               |               |                    |                   |                    |                   |                     |               |               |                    |                   |                    |                    |                |                   |
         // generic instances
-        // TODO: Change to ScalarPerVector = 1 when inner_product<half_t, half_t, float> will be supported
         DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK<           2,    F16,     F16,  DsDatatype,     F16,     F32, InLayout, WeiLayout,  DsLayout, OutLayout, PassThrough,  PassThrough,  CDEElementOp,       ConvSpec,    GemmMNKPadding,     8,    16,     4,     2,  1,          1,          2,      1,       S<4, 2>,       S<1, 1>,      S<2, 1, 2, 1>,        S<1, 1, 8, 1>,  S<1, 2, 0, 3>,  S<1, 2, 0, 3>,       S<1, 1, 1, 1>,      S<1, 2, 0, 3>,       S<1, 1, 1, 1>,      S<1, 1, 1, 1>,        S<2, 1, 4, 1>,  S<1, 2, 0, 3>,  S<1, 2, 0, 3>,       S<1, 1, 1, 1>,      S<1, 2, 0, 3>,       S<1, 1, 1, 1>, S<0, 1, 2, 3, 4, 5>,               5,                 1>,
         DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK<           2,    F16,     F16,  DsDatatype,     F16,     F32, InLayout, WeiLayout,  DsLayout, OutLayout, PassThrough,  PassThrough,  CDEElementOp,       ConvSpec,    GemmMNKPadding,   256,   128,   128,    16,  1,          4,          4,      1,       S<8, 2>,       S<8, 2>,      S<8, 1, 1, 1>,      S<2, 1, 128, 1>,  S<1, 2, 0, 3>,  S<1, 2, 0, 3>,       S<1, 1, 1, 1>,      S<1, 2, 0, 3>,       S<1, 1, 1, 1>,      S<8, 1, 1, 1>,      S<2, 1, 128, 1>,  S<1, 2, 0, 3>,  S<1, 2, 0, 3>,       S<1, 1, 1, 1>,      S<1, 2, 0, 3>,       S<1, 1, 1, 1>, S<0, 1, 2, 3, 4, 5>,               5,                 1>,
 

test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp

@@ -33,8 +33,8 @@ class TestGroupedConvndBwdWeight : public ::testing::Test
 
     bool skip_case(const ck::utils::conv::ConvParam& params, const ck::index_t split_k)
     {
-        // K or C are odd is supported only by DL kernel (only applies to fp16)
-        // DL kernel is only supported for split_k=1
+        // Odd K or C values are supported only by DL kernel (only applies to fp16)
+        // DL kernel currently supports only `split_k=1`
         if constexpr(std::is_same_v<InDataType, ck::half_t>)
         {
             if(split_k != 1 && (params.K_ % 2 != 0 || params.C_ % 2 != 0))