Stylistic improvements for grouped convolution code

Remove unnecessary ignoring Update test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp

Stylistic improvements for grouped convolution code
Remove unnecessary ignoring Update test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp
bc2d0583 · Bartlomiej Kocot · Bartłomiej Kocot · f9d0eddb · bc2d0583 · bc2d0583
Commit bc2d0583 authored Sep 13, 2023 by Bartlomiej Kocot Committed by Bartłomiej Kocot Sep 15, 2023
3 changed files
--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp
+++ b/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
 }


--- a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_dl_instance.hpp
+++ b/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>,


--- a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight.cpp
+++ b/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))