Conv3D FWD BWD WRW fp16 fp32 client examples

client_example/11_grouped_conv_bwd_weight/CMakeLists.txt

@@ -1,2 +1,7 @@
-add_executable(client_grouped_conv2d_bwd_weight grouped_conv2d_bwd_weight.cpp)
-target_link_libraries(client_grouped_conv2d_bwd_weight PRIVATE composable_kernel::device_operations)
+add_executable(client_grouped_conv2d_bwd_weight_fp16 grouped_conv2d_bwd_weight_fp16.cpp)
+add_executable(client_grouped_conv3d_bwd_weight_fp16 grouped_conv3d_bwd_weight_fp16.cpp)
+add_executable(client_grouped_conv3d_bwd_weight_fp32 grouped_conv3d_bwd_weight_fp32.cpp)
+
+target_link_libraries(client_grouped_conv2d_bwd_weight_fp16 PRIVATE composable_kernel::device_operations)
+target_link_libraries(client_grouped_conv3d_bwd_weight_fp16 PRIVATE composable_kernel::device_operations)
+target_link_libraries(client_grouped_conv3d_bwd_weight_fp32 PRIVATE composable_kernel::device_operations)

client_example/11_grouped_conv_bwd_weight/grouped_conv2d_bwd_weight.cpp → client_example/11_grouped_conv_bwd_weight/common.hpp

@@ -13,27 +13,8 @@
 #include "ck/tensor_operation/gpu/device/device_conv_fwd.hpp"
 #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
 
-using InDataType  = ck::half_t;
-using WeiDataType = ck::half_t;
-using OutDataType = ck::half_t;
-
-using InLayout    = ck::tensor_layout::convolution::GNHWC;
-using WeiLayout   = ck::tensor_layout::convolution::GKYXC;
-using OutLayout   = ck::tensor_layout::convolution::GNHWK;
 using PassThrough = ck::tensor_operation::element_wise::PassThrough;
 
-static constexpr ck::index_t NumDimSpatial = 2;
-static constexpr ck::index_t G             = 32;
-static constexpr ck::index_t N             = 256;
-static constexpr ck::index_t K             = 192;
-static constexpr ck::index_t C             = 192;
-static constexpr ck::index_t Y             = 3;
-static constexpr ck::index_t X             = 3;
-static constexpr ck::index_t Hi            = 28;
-static constexpr ck::index_t Wi            = 28;
-static constexpr ck::index_t Ho            = 28;
-static constexpr ck::index_t Wo            = 28;
-
 struct SimpleDeviceMem
 {
     SimpleDeviceMem() = delete;
@@ -50,22 +31,93 @@ struct SimpleDeviceMem
     void* p_mem_;
 };
 
-int main()
+template <ck::index_t NumDimSpatial>
+std::size_t GetFlops(ck::index_t G,
+                     ck::index_t N,
+                     ck::index_t K,
+                     ck::index_t C,
+                     const std::array<ck::index_t, NumDimSpatial>& output_spatial_lengths,
+                     const std::array<ck::index_t, NumDimSpatial>& filter_spatial_lengths)
 {
-    std::array<ck::index_t, NumDimSpatial> input_spatial_lengths{Hi, Wi};
-    std::array<ck::index_t, NumDimSpatial> filter_spatial_lengths{Y, X};
-    std::array<ck::index_t, NumDimSpatial> output_spatial_lengths{Ho, Wo};
+    // 2 * G * N * K * C * <output spatial lengths product> * <filter spatial lengths product>
+    return static_cast<std::size_t>(2) * G * N * K * C *
+           std::accumulate(std::begin(output_spatial_lengths),
+                           std::end(output_spatial_lengths),
+                           static_cast<std::size_t>(1),
+                           std::multiplies<>()) *
+           std::accumulate(std::begin(filter_spatial_lengths),
+                           std::end(filter_spatial_lengths),
+                           static_cast<std::size_t>(1),
+                           std::multiplies<>());
+}
 
-    std::array<ck::index_t, NumDimSpatial> conv_filter_strides{1, 1};
-    std::array<ck::index_t, NumDimSpatial> conv_filter_dilations{1, 1};
-    std::array<ck::index_t, NumDimSpatial> input_left_pads{1, 1};
-    std::array<ck::index_t, NumDimSpatial> input_right_pads{1, 1};
+template <typename InDataType, ck::index_t NumDimSpatial>
+std::size_t GetInputByte(ck::index_t G,
+                         ck::index_t N,
+                         ck::index_t C,
+                         const std::array<ck::index_t, NumDimSpatial>& input_spatial_lengths)
+{
+    // sizeof(InDataType) * (G * N * C * <input spatial lengths product>) +
+    return sizeof(InDataType) * (G * N * C *
+                                 std::accumulate(std::begin(input_spatial_lengths),
+                                                 std::end(input_spatial_lengths),
+                                                 static_cast<std::size_t>(1),
+                                                 std::multiplies<>()));
+}
 
-    ck::index_t split_k = 2;
+template <typename WeiDataType, ck::index_t NumDimSpatial>
+std::size_t GetWeightByte(ck::index_t G,
+                          ck::index_t K,
+                          ck::index_t C,
+                          const std::array<ck::index_t, NumDimSpatial>& filter_spatial_lengths)
+{
+    // sizeof(WeiDataType) * (G * K * C * <filter spatial lengths product>) +
+    return sizeof(WeiDataType) * (G * K * C *
+                                  std::accumulate(std::begin(filter_spatial_lengths),
+                                                  std::end(filter_spatial_lengths),
+                                                  static_cast<std::size_t>(1),
+                                                  std::multiplies<>()));
+}
 
-    SimpleDeviceMem in(sizeof(InDataType) * G * N * Hi * Wi * C);
-    SimpleDeviceMem wei(sizeof(WeiDataType) * G * K * Y * X * C);
-    SimpleDeviceMem out(sizeof(OutDataType) * G * N * Ho * Wo * K);
+template <typename OutDataType, ck::index_t NumDimSpatial>
+std::size_t GetOutputByte(ck::index_t G,
+                          ck::index_t N,
+                          ck::index_t K,
+                          const std::array<ck::index_t, NumDimSpatial>& output_spatial_lengths)
+{
+    // sizeof(OutDataType) * (G * N * K * <output spatial lengths product>);
+    return sizeof(OutDataType) * (G * N * K *
+                                  std::accumulate(std::begin(output_spatial_lengths),
+                                                  std::end(output_spatial_lengths),
+                                                  static_cast<std::size_t>(1),
+                                                  std::multiplies<std::size_t>()));
+}
+
+template <ck::index_t NumDimSpatial,
+          typename InDataType,
+          typename WeiDataType,
+          typename OutDataType,
+          typename InLayout,
+          typename WeiLayout,
+          typename OutLayout>
+bool run_grouped_conv_bwd_weight(
+    ck::index_t G,
+    ck::index_t N,
+    ck::index_t K,
+    ck::index_t C,
+    const std::array<ck::index_t, NumDimSpatial>& input_spatial_lengths,
+    const std::array<ck::index_t, NumDimSpatial>& filter_spatial_lengths,
+    const std::array<ck::index_t, NumDimSpatial>& output_spatial_lengths,
+    const std::array<ck::index_t, NumDimSpatial>& conv_filter_strides,
+    const std::array<ck::index_t, NumDimSpatial>& conv_filter_dilations,
+    const std::array<ck::index_t, NumDimSpatial>& input_left_pads,
+    const std::array<ck::index_t, NumDimSpatial>& input_right_pads)
+{
+
+    ck::index_t split_k = 2;
+    SimpleDeviceMem in(GetInputByte<InDataType, NumDimSpatial>(G, N, C, input_spatial_lengths));
+    SimpleDeviceMem wei(GetWeightByte<WeiDataType, NumDimSpatial>(G, K, C, filter_spatial_lengths));
+    SimpleDeviceMem out(GetOutputByte<OutDataType, NumDimSpatial>(G, N, K, output_spatial_lengths));
 
     using DeviceOp = ck::tensor_operation::device::DeviceGroupedConvBwdWeight<NumDimSpatial,
                                                                               InLayout,
@@ -120,10 +172,12 @@ int main()
         {
             float avg_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, true});
 
-            std::size_t flop      = std::size_t(2) * G * N * K * C * Ho * Wo * Y * X;
-            std::size_t num_bytes = sizeof(InDataType) * G * N * Hi * Wi * C +
-                                    sizeof(WeiDataType) * G * K * Y * X * C +
-                                    sizeof(OutDataType) * G * N * Ho * Wo * K;
+            std::size_t flop =
+                GetFlops<NumDimSpatial>(G, N, K, C, output_spatial_lengths, filter_spatial_lengths);
+            std::size_t num_bytes =
+                GetInputByte<InDataType, NumDimSpatial>(G, N, C, input_spatial_lengths) +
+                GetWeightByte<WeiDataType, NumDimSpatial>(G, K, C, filter_spatial_lengths) +
+                GetOutputByte<OutDataType, NumDimSpatial>(G, N, K, output_spatial_lengths);
 
             float tflops     = static_cast<float>(flop) / 1.E9 / avg_time;
             float gb_per_sec = num_bytes / 1.E6 / avg_time;
@@ -149,7 +203,7 @@ int main()
     if(best_op_id < 0)
     {
         std::cerr << "no suitable instance" << std::endl;
-        return EXIT_FAILURE;
+        return false;
     }
 
     std::cout << "Best Perf: " << std::setw(10) << best_avg_time << " ms, " << best_tflops
@@ -187,4 +241,6 @@ int main()
 
         std::cout << "Done" << std::endl;
     }
+
+    return true;
 }

client_example/11_grouped_conv_bwd_weight/grouped_conv2d_bwd_weight_fp16.cpp

@@ -0,0 +1,41 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "common.hpp"
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+
+using InDataType  = ck::half_t;
+using WeiDataType = ck::half_t;
+using OutDataType = ck::half_t;
+
+using InLayout  = ck::tensor_layout::convolution::GNHWC;
+using WeiLayout = ck::tensor_layout::convolution::GKYXC;
+using OutLayout = ck::tensor_layout::convolution::GNHWK;
+
+static constexpr ck::index_t NumDimSpatial = 2;
+static constexpr ck::index_t G             = 32;
+static constexpr ck::index_t N             = 256;
+static constexpr ck::index_t K             = 192;
+static constexpr ck::index_t C             = 192;
+static constexpr ck::index_t Y             = 3;
+static constexpr ck::index_t X             = 3;
+static constexpr ck::index_t Hi            = 28;
+static constexpr ck::index_t Wi            = 28;
+static constexpr ck::index_t Ho            = 28;
+static constexpr ck::index_t Wo            = 28;
+
+int main()
+{
+    return run_grouped_conv_bwd_weight<NumDimSpatial,
+                                       InDataType,
+                                       WeiDataType,
+                                       OutDataType,
+                                       InLayout,
+                                       WeiLayout,
+                                       OutLayout>(
+               G, N, K, C, {Hi, Wi}, {Y, X}, {Ho, Wo}, {1, 1}, {1, 1}, {1, 1}, {1, 1})
+               ? EXIT_SUCCESS
+               : EXIT_FAILURE;
+}

client_example/11_grouped_conv_bwd_weight/grouped_conv3d_bwd_weight_fp16.cpp

@@ -0,0 +1,53 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "common.hpp"
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+
+using InDataType  = ck::half_t;
+using WeiDataType = ck::half_t;
+using OutDataType = ck::half_t;
+
+using InLayout  = ck::tensor_layout::convolution::GNDHWC;
+using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
+using OutLayout = ck::tensor_layout::convolution::GNDHWK;
+
+static constexpr ck::index_t NumDimSpatial = 3;
+static constexpr ck::index_t G             = 8;
+static constexpr ck::index_t N             = 64;
+static constexpr ck::index_t K             = 128;
+static constexpr ck::index_t C             = 128;
+static constexpr ck::index_t Z             = 3;
+static constexpr ck::index_t Y             = 3;
+static constexpr ck::index_t X             = 3;
+static constexpr ck::index_t Di            = 28;
+static constexpr ck::index_t Hi            = 28;
+static constexpr ck::index_t Wi            = 3;
+static constexpr ck::index_t Do            = 28;
+static constexpr ck::index_t Ho            = 28;
+static constexpr ck::index_t Wo            = 3;
+
+int main()
+{
+    return run_grouped_conv_bwd_weight<NumDimSpatial,
+                                       InDataType,
+                                       WeiDataType,
+                                       OutDataType,
+                                       InLayout,
+                                       WeiLayout,
+                                       OutLayout>(G,
+                                                  N,
+                                                  K,
+                                                  C,
+                                                  {Di, Hi, Wi},
+                                                  {Z, Y, X},
+                                                  {Do, Ho, Wo},
+                                                  {1, 1, 1},
+                                                  {1, 1, 1},
+                                                  {1, 1, 1},
+                                                  {1, 1, 1})
+               ? EXIT_SUCCESS
+               : EXIT_FAILURE;
+}

client_example/11_grouped_conv_bwd_weight/grouped_conv3d_bwd_weight_fp32.cpp

@@ -0,0 +1,53 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "common.hpp"
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+
+using InDataType  = float;
+using WeiDataType = float;
+using OutDataType = float;
+
+using InLayout  = ck::tensor_layout::convolution::GNDHWC;
+using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
+using OutLayout = ck::tensor_layout::convolution::GNDHWK;
+
+static constexpr ck::index_t NumDimSpatial = 3;
+static constexpr ck::index_t G             = 8;
+static constexpr ck::index_t N             = 64;
+static constexpr ck::index_t K             = 128;
+static constexpr ck::index_t C             = 128;
+static constexpr ck::index_t Z             = 3;
+static constexpr ck::index_t Y             = 3;
+static constexpr ck::index_t X             = 3;
+static constexpr ck::index_t Di            = 28;
+static constexpr ck::index_t Hi            = 28;
+static constexpr ck::index_t Wi            = 3;
+static constexpr ck::index_t Do            = 28;
+static constexpr ck::index_t Ho            = 28;
+static constexpr ck::index_t Wo            = 3;
+
+int main()
+{
+    return run_grouped_conv_bwd_weight<NumDimSpatial,
+                                       InDataType,
+                                       WeiDataType,
+                                       OutDataType,
+                                       InLayout,
+                                       WeiLayout,
+                                       OutLayout>(G,
+                                                  N,
+                                                  K,
+                                                  C,
+                                                  {Di, Hi, Wi},
+                                                  {Z, Y, X},
+                                                  {Do, Ho, Wo},
+                                                  {1, 1, 1},
+                                                  {1, 1, 1},
+                                                  {1, 1, 1},
+                                                  {1, 1, 1})
+               ? EXIT_SUCCESS
+               : EXIT_FAILURE;
+}

client_example/15_convnd_bwd_data/CMakeLists.txt

@@ -0,0 +1,5 @@
+add_executable(client_conv3d_bwd_data_fp16 conv3d_bwd_data_fp16.cpp)
+add_executable(client_conv3d_bwd_data_fp32 conv3d_bwd_data_fp32.cpp)
+
+target_link_libraries(client_conv3d_bwd_data_fp16 PRIVATE composable_kernel::device_operations)
+target_link_libraries(client_conv3d_bwd_data_fp32 PRIVATE composable_kernel::device_operations)