add pass by point

composable_kernel/include/tensor_operation/gridwise_gemm_dlops_v3.hpp

@@ -13,6 +13,45 @@
 namespace ck {
 
 #if CK_EXPERIMENTAL_PASS_TENSOR_DESCRIPTOR_BY_VALUE
+template <typename GridwiseGemm,
+          typename FloatAB,
+          typename FloatC,
+          typename AGridDesc_E0_E1_K0_K1_E2,
+          typename BGridDesc_E0_E1_N_H0_H1_H2_W0_W1_W2_E2,
+          typename CGridDesc_K0_K1_N_H0_H1_H2_W0_W1_W2,
+          typename CBlockIdToBlockClusterAdaptor_K_N_H_W,
+          bool HasMainE0BlockLoop>
+__global__ void
+#if CK_USE_LAUNCH_BOUNDS
+    __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
+#endif
+        kernel_gemm_dlops_v3(
+            const FloatAB* __restrict__ p_a_grid,
+            const FloatAB* __restrict__ p_b_grid,
+            const FloatC* __restrict__ p_bias_grid,
+            FloatC* __restrict__ p_c_grid,
+            const AGridDesc_E0_E1_K0_K1_E2 a_e0_e1_k0_k1_e2_grid_desc,
+            const BGridDesc_E0_E1_N_H0_H1_H2_W0_W1_W2_E2 b_e0_e1_n_h0_h1_h2_w0_w1_w2_e2_grid_desc,
+            const CGridDesc_K0_K1_N_H0_H1_H2_W0_W1_W2 c_k0_k1_n_h0_h1_h2_w0_w1_w2_grid_desc,
+            const CBlockIdToBlockClusterAdaptor_K_N_H_W c_blockid_to_k_n_h_w_block_cluster_adaptor)
+{
+    constexpr index_t shared_block_size =
+        GridwiseGemm::GetSharedMemoryNumberOfByte() / sizeof(FloatAB);
+
+    __shared__ FloatAB p_shared_block[shared_block_size];
+
+    GridwiseGemm::ConvBiasActiv(p_a_grid,
+                                p_b_grid,
+                                p_bias_grid,
+                                p_c_grid,
+                                p_shared_block,
+                                a_e0_e1_k0_k1_e2_grid_desc,
+                                b_e0_e1_n_h0_h1_h2_w0_w1_w2_e2_grid_desc,
+                                c_k0_k1_n_h0_h1_h2_w0_w1_w2_grid_desc,
+                                c_blockid_to_k_n_h_w_block_cluster_adaptor,
+                                integral_constant<bool, HasMainE0BlockLoop>{});
+}
+
 template <typename GridwiseGemm,
           typename FloatAB,
           typename FloatC,
@@ -26,7 +65,49 @@ __global__ void
 #if CK_USE_LAUNCH_BOUNDS
     __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
 #endif
-        kernel_gemm_dlops_v2_add(
+        kernel_gemm_dlops_v3_resize_add(
+            const FloatAB* __restrict__ p_a_grid,
+            const FloatAB* __restrict__ p_b_grid,
+            const FloatC* __restrict__ p_bias_grid,
+            FloatC* __restrict__ p_d_grid,
+            const AGridDesc_E0_E1_K0_K1_E2 a_e0_e1_k0_k1_e2_grid_desc,
+            const BGridDesc_E0_E1_N_H0_H1_H2_W0_W1_W2_E2 b_e0_e1_n_h0_h1_h2_w0_w1_w2_e2_grid_desc,
+            const CGridDesc_K0_K1_N_H0_H1_H2_W0_W1_W2 c_k0_k1_n_h0_h1_h2_w0_w1_w2_grid_desc,
+            const DGridDesc_K0_K1_N_H0_H1_Hx_W0_W1_Wx d_k0_k1_n_h0_h1_hx_w0_w1_wx_grid_desc,
+            const CBlockIdToBlockClusterAdaptor_K_N_H_W c_blockid_to_k_n_h_w_block_cluster_adaptor)
+{
+    constexpr index_t shared_block_size =
+        GridwiseGemm::GetSharedMemoryNumberOfByte() / sizeof(FloatAB);
+
+    __shared__ FloatAB p_shared_block[shared_block_size];
+
+    GridwiseGemm::ConvBiasActivResizeAddRun(p_a_grid,
+                                            p_b_grid,
+                                            p_bias_grid,
+                                            p_d_grid,
+                                            p_shared_block,
+                                            a_e0_e1_k0_k1_e2_grid_desc,
+                                            b_e0_e1_n_h0_h1_h2_w0_w1_w2_e2_grid_desc,
+                                            c_k0_k1_n_h0_h1_h2_w0_w1_w2_grid_desc,
+                                            d_k0_k1_n_h0_h1_hx_w0_w1_wx_grid_desc,
+                                            c_blockid_to_k_n_h_w_block_cluster_adaptor,
+                                            integral_constant<bool, HasMainE0BlockLoop>{});
+}
+
+template <typename GridwiseGemm,
+          typename FloatAB,
+          typename FloatC,
+          typename AGridDesc_E0_E1_K0_K1_E2,
+          typename BGridDesc_E0_E1_N_H0_H1_H2_W0_W1_W2_E2,
+          typename CGridDesc_K0_K1_N_H0_H1_H2_W0_W1_W2,
+          typename DGridDesc_K0_K1_N_H0_H1_Hx_W0_W1_Wx,
+          typename CBlockIdToBlockClusterAdaptor_K_N_H_W,
+          bool HasMainE0BlockLoop>
+__global__ void
+#if CK_USE_LAUNCH_BOUNDS
+    __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, CK_MIN_BLOCK_PER_CU)
+#endif
+        kernel_gemm_dlops_v3_maxpool(
             const FloatAB* __restrict__ p_a_grid,
             const FloatAB* __restrict__ p_b_grid,
             const FloatC* __restrict__ p_bias_grid,
@@ -43,18 +124,18 @@ __global__ void
 
     __shared__ FloatAB p_shared_block[shared_block_size];
 
-    GridwiseGemm::Run(p_a_grid,
-                      p_b_grid,
-                      p_bias_grid,
-                      p_c_grid,
-                      p_d_grid,
-                      p_shared_block,
-                      a_e0_e1_k0_k1_e2_grid_desc,
-                      b_e0_e1_n_h0_h1_h2_w0_w1_w2_e2_grid_desc,
-                      c_k0_k1_n_h0_h1_h2_w0_w1_w2_grid_desc,
-                      d_k0_k1_n_h0_h1_hx_w0_w1_wx_grid_desc,
-                      c_blockid_to_k_n_h_w_block_cluster_adaptor,
-                      integral_constant<bool, HasMainE0BlockLoop>{});
+    GridwiseGemm::ConvBiasActivMaxpoolRun(p_a_grid,
+                                          p_b_grid,
+                                          p_bias_grid,
+                                          p_c_grid,
+                                          p_d_grid,
+                                          p_shared_block,
+                                          a_e0_e1_k0_k1_e2_grid_desc,
+                                          b_e0_e1_n_h0_h1_h2_w0_w1_w2_e2_grid_desc,
+                                          c_k0_k1_n_h0_h1_h2_w0_w1_w2_grid_desc,
+                                          d_k0_k1_n_h0_h1_hx_w0_w1_wx_grid_desc,
+                                          c_blockid_to_k_n_h_w_block_cluster_adaptor,
+                                          integral_constant<bool, HasMainE0BlockLoop>{});
 }
 #elif CK_EXPERIMENTAL_PASS_TENSOR_DESCRIPTOR_BY_VOID_POINTER
 // pass tensor descriptor by CONSTANT void pointer

composable_kernel/include/utility/config.hpp

@@ -90,8 +90,8 @@
 #endif
 
 // pass tensor descriptor by value or void*
-#define CK_EXPERIMENTAL_PASS_TENSOR_DESCRIPTOR_BY_VALUE 0
-#define CK_EXPERIMENTAL_PASS_TENSOR_DESCRIPTOR_BY_VOID_POINTER 1
+#define CK_EXPERIMENTAL_PASS_TENSOR_DESCRIPTOR_BY_VALUE 1
+#define CK_EXPERIMENTAL_PASS_TENSOR_DESCRIPTOR_BY_VOID_POINTER 0
 #define CK_EXPERIMENTAL_STATIC_TENSOR_DESCRIPTOR 0
 
 // merge transformation use magic number division

host/driver_offline/include/driver_convolution_add_forward_implicit_gemm_v5r1_dlops_nc0hwc1_kc0yxc1_nk0hwk1.hpp

@@ -373,7 +373,7 @@ struct DriverDynamicConvolutionForwardImplicitGemmDlops_v5r1_nc0hwc1_kc0yxc1_nk0
 
         if(has_main_e0_block_loop)
         {
-            const auto kernel = kernel_gemm_dlops_v2_resize_add<
+            const auto kernel = kernel_gemm_dlops_v3_resize_add<
                 GridwiseGemm,
                 FloatAB,
                 FloatC,
@@ -401,7 +401,7 @@ struct DriverDynamicConvolutionForwardImplicitGemmDlops_v5r1_nc0hwc1_kc0yxc1_nk0
         }
         else
         {
-            const auto kernel = kernel_gemm_dlops_v2_resize_add<
+            const auto kernel = kernel_gemm_dlops_v3_resize_add<
                 GridwiseGemm,
                 FloatAB,
                 FloatC,

host/driver_offline/include/driver_convolution_forward_implicit_gemm_v5r1_dlops_nc0hwc1_kc0yxc1_nk0hwk1.hpp

@@ -330,7 +330,7 @@ struct DriverDynamicConvolutionForwardImplicitGemmDlops_v5r1_nc0hwc1_kc0yxc1_nk0
         if(has_main_e0_block_loop)
         {
             const auto kernel =
-                kernel_gemm_dlops_v2<GridwiseGemm,
+                kernel_gemm_dlops_v3<GridwiseGemm,
                                      FloatAB,
                                      FloatC,
                                      remove_reference_t<AGridDesc_E0_E1_K0_K1_E2>,
@@ -356,7 +356,7 @@ struct DriverDynamicConvolutionForwardImplicitGemmDlops_v5r1_nc0hwc1_kc0yxc1_nk0
         else
         {
             const auto kernel =
-                kernel_gemm_dlops_v2<GridwiseGemm,
+                kernel_gemm_dlops_v3<GridwiseGemm,
                                      FloatAB,
                                      FloatC,
                                      remove_reference_t<AGridDesc_E0_E1_K0_K1_E2>,