ckProfiler and device-level XDL GEMM operator

CMakeLists.txt

@@ -40,6 +40,7 @@ message(STATUS "Build with HIP ${hip_VERSION}")
 
 ## half
 #find_path(HALF_INCLUDE_DIR half.hpp)
+set(HALF_INCLUDE_DIR "${PROJECT_SOURCE_DIR}/external/half/include")
 message("HALF_INCLUDE_DIR: ${HALF_INCLUDE_DIR}")
 
 # CMAKE_CXX_FLAGS
@@ -185,6 +186,7 @@ enable_cppcheck(
         composable_kernel/src/kernel_wrapper
     INCLUDE
         host/host_tensor/include
+        host/device/include
         host/solver/include
         host/driver_offline/include
         composable_kernel/include/*
@@ -196,3 +198,5 @@ enable_cppcheck(
 )
 
 add_subdirectory(host)
+add_subdirectory(example)
+add_subdirectory(profiler)

README.md

@@ -1,177 +1 @@
-# How to build and run
 
-# Docker
-```
-docker run                                                                   \
--it                                                                          \
---rm                                                                         \
---privileged                                                                 \
---group-add sudo                                                             \
--w /root/workspace                                                           \
--v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace                                \
-rocm/tensorflow:rocm4.2-tf2.4-dev                                            \
-/bin/bash
-```
-
-# Install Boost for online compilation
-https://www.boost.org/doc/libs/1_66_0/more/getting_started/unix-variants.html#easy-build-and-install
-
-
-# Build
-Add path of Boost
-```
- export LD_LIBRARY_PATH=/usr/local/lib:$LD_LIBRARY_PATH
-```
-
-```
-mkdir build && cd build
-```
-
-cmake cmd. Need to Specify target ID, example below is gfx908
-```
-cmake                                                                                                                              \
--D CMAKE_BUILD_TYPE=Release                                                                                                                    \
--D CMAKE_CXX_FLAGS="-DCK_AMD_GPU_GFX908 -O3 --amdgpu-target=gfx908 -mllvm --amdgpu-spill-vgpr-to-agpr=0 -gline-tables-only -save-temps=$PWD"   \
--D HIP_ONLINE_COMPILER_FLAGS="-DCK_AMD_GPU_GFX908"                                                                                             \
--D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                                                                      \
--D CMAKE_PREFIX_PATH=/opt/rocm                                                                                                                 \
--D CMAKE_VERBOSE_MAKEFILE:BOOL=ON                                                                                                              \
-..
-```
-
-Build drivers:   \
-``conv_fwd_driver_offline`` is (offline compilation) driver for forward convolution,  \
-``conv_bwd_driver_offline`` is (offline compilation) driver for backward-data convolution  \
-``conv_fwd_driver_online`` is (online compilation) driver for forward convolution
-```
- make -j conv_fwd_driver_offline
- make -j conv_bwd_driver_offline
- make -j conv_fwd_driver_online
-```
-
-# Run
-* layout: 0 = NCHW; 1 = NHWC
-* algo: algorithm
-* verify: 0 = no verification; 1 = do verification
-* init: 0 ~ 5. initialization method
-* log: 0 = no log; 1 = do log
-* repeat: number of time kernel being launched
-```
-######################################################## layout  algo  verify  init  log  repeat  N__ K___ C___ Y X Hi_ Wi__ Strides Dilations LeftPads RightPads
- ./host/driver_offline/conv_fwd_driver_offline                0     4       0     0    0       1  128  256  192 3 3  71   71     2 2       1 1      1 1       1 1
- ./host/driver_offline/conv_fwd_driver_offline                0     4       0     0    0       1  256 1024  256 3 3  14   14     1 1       1 1      1 1       1 1
- ./host/driver_offline/conv_fwd_driver_offline                1     5       0     0    0       1  128  256  192 3 3  71   71     2 2       1 1      1 1       1 1
- ./host/driver_offline/conv_fwd_driver_offline                1     5       0     0    0       1  256 1024  256 3 3  14   14     1 1       1 1      1 1       1 1
- ./host/driver_offline/conv_bwd_driver_offline                1     5       0     0    0       1  256  256 1024 3 3  14   14     1 1       1 1      1 1       1 1
-```
-
-# Result
-Forward convoltuion, FP16, NCHW
-```
-./host/driver_offline/conv_fwd_driver_offline                0     4       0     0    0       1  128  256  192 3 3  71   71     2 2       1 1      1 1       1 1
-
-layout: 0
-in: dim 4, lengths {128, 192, 71, 71}, strides {967872, 5041, 71, 1}
-wei: dim 4, lengths {256, 192, 3, 3}, strides {1728, 9, 3, 1}
-out: dim 4, lengths {128, 256, 36, 36}, strides {331776, 1296, 36, 1}
-InLeftPads size 2, {1, 1, }
-InRightPads size 2, {1, 1, }
-ConvStrides size 2, {2, 2, }
-ConvDilations size 2, {1, 1, }
-device_convolution_forward_implicit_gemm_v4r4r2_xdlops_nchw_kcyx_nkhw
-a_k0_m_k1_grid_desc{216, 256, 8}
-b_k0_n_k1_grid_desc{216, 165888, 8}
-c_m_n_grid_desc{ 256, 165888}
-launch_and_time_kernel: grid_dim {1296, 1, 1}, block_dim {256, 1, 1}
-Warm up
-Start running 1 times...
-Average time : 1.4155 ms, 103.686 TFlop/s
-```
-
-Forward convoltuion, FP16, NCHW
-```
- ./host/driver_offline/conv_fwd_driver_offline                0     4       0     0    0       1  256 1024  256 3 3  14   14     1 1       1 1      1 1       1 1
-
- layout: 0
-in: dim 4, lengths {256, 256, 14, 14}, strides {50176, 196, 14, 1}
-wei: dim 4, lengths {1024, 256, 3, 3}, strides {2304, 9, 3, 1}
-out: dim 4, lengths {256, 1024, 14, 14}, strides {200704, 196, 14, 1}
-InLeftPads size 2, {1, 1, }
-InRightPads size 2, {1, 1, }
-ConvStrides size 2, {1, 1, }
-ConvDilations size 2, {1, 1, }
-device_convolution_forward_implicit_gemm_v4r4r2_xdlops_nchw_kcyx_nkhw
-a_k0_m_k1_grid_desc{288, 1024, 8}
-b_k0_n_k1_grid_desc{288, 50176, 8}
-c_m_n_grid_desc{ 1024, 50176}
-launch_and_time_kernel: grid_dim {1568, 1, 1}, block_dim {256, 1, 1}
-Warm up
-Start running 1 times...
-Average time : 2.21357 ms, 106.959 TFlop/s
- ```
-
- Forward convolution, FP16, NHWC
- ```
- ./host/driver_offline/conv_fwd_driver_offline                1     5       0     0    0       1  128  256  192 3 3  71   71     2 2       1 1      1 1       1 1
-
- layout: 1
-in: dim 4, lengths {128, 71, 71, 192}, strides {967872, 13632, 192, 1}
-wei: dim 4, lengths {256, 3, 3, 192}, strides {1728, 576, 192, 1}
-out: dim 4, lengths {128, 36, 36, 256}, strides {331776, 9216, 256, 1}
-InLeftPads size 2, {1, 1, }
-InRightPads size 2, {1, 1, }
-ConvStrides size 2, {2, 2, }
-ConvDilations size 2, {1, 1, }
-device_convolution_forward_implicit_gemm_v4r4r4_xdlops_nhwc_kyxc_nhwk
-a_k0_m_k1_grid_desc{216, 165888, 8}
-b_k0_n_k1_grid_desc{216, 256, 8}
-c_m_n_grid_desc{ 165888, 256}
-launch_and_time_kernel: grid_dim {1296, 1, 1}, block_dim {256, 1, 1}
-Warm up
-Start running 1 times...
-Average time : 1.12014 ms, 131.025 TFlop/s
- ```
-
- Forward convolution, FP16, NHWC
- ```
- ./host/driver_offline/conv_fwd_driver_offline                1     5       0     0    0       1  256 1024  256 3 3  14   14     1 1       1 1      1 1       1 1
-
- layout: 1
-in: dim 4, lengths {256, 14, 14, 256}, strides {50176, 3584, 256, 1}
-wei: dim 4, lengths {1024, 3, 3, 256}, strides {2304, 768, 256, 1}
-out: dim 4, lengths {256, 14, 14, 1024}, strides {200704, 14336, 1024, 1}
-InLeftPads size 2, {1, 1, }
-InRightPads size 2, {1, 1, }
-ConvStrides size 2, {1, 1, }
-ConvDilations size 2, {1, 1, }
-device_convolution_forward_implicit_gemm_v4r4r4_xdlops_nhwc_kyxc_nhwk
-a_k0_m_k1_grid_desc{288, 50176, 8}
-b_k0_n_k1_grid_desc{288, 1024, 8}
-c_m_n_grid_desc{ 50176, 1024}
-launch_and_time_kernel: grid_dim {1568, 1, 1}, block_dim {256, 1, 1}
-Warm up
-Start running 1 times...
-Average time : 1.86877 ms, 126.693 TFlop/s
- ```
-
- Backward data convolution, FP16, NHWC
- ```
- ./host/driver_offline/conv_bwd_driver_offline       1     1       0     3    0       1  256  256 1024 3 3  14   14     1 1       1 1      1 1       1 1
-
- layout: 1
-in: dim 4, lengths {256, 14, 14, 1024}, strides {200704, 14336, 1024, 1}
-wei: dim 4, lengths {256, 3, 3, 1024}, strides {9216, 3072, 1024, 1}
-out: dim 4, lengths {256, 14, 14, 256}, strides {50176, 3584, 256, 1}
-InLeftPads size 2, {1, 1, }
-InRightPads size 2, {1, 1, }
-ConvStrides size 2, {1, 1, }
-ConvDilations size 2, {1, 1, }
-device_convolution_backward_data_implicit_gemm_v4r1r2_xdlops_nhwc_kyxc_nhwk
-a_k0_m_k1_grid_desc{288, 50176, 8}
-b_k0_n_k1_grid_desc{288, 1024, 8}
-c_m_n_grid_desc{ 50176, 1024}
-launch_and_time_kernel: grid_dim {1568, 1, 1}, block_dim {256, 1, 1}
-Warm up
-Start running 1 times...
-Average time : 2.22461 ms, 106.428 TFlop/s
-```

composable_kernel/include/problem_transform/transform_forward_convolution_into_gemm_v4r4r4_nhwc_kyxc_nhwk.hpp

@@ -21,8 +21,7 @@ template <typename... In,
           typename InLeftPads,
           typename InRightPads,
           index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_forward_convolution_into_gemm_v4r4r4_nhwc_kyxc_nhwk_pad(
+__host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4r4_nhwc_kyxc_nhwk(
     const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
     const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
     const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,

composable_kernel/include/tensor_operation/blockwise_gemm_xdlops.hpp

@@ -124,7 +124,7 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
                       "wrong!");
     }
 
-    __host__ __device__ static constexpr auto GetCM0N0M1N1M2M3M4N2ThreadDescriptor()
+    __host__ __device__ static constexpr auto GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2()
     {
         constexpr auto c_m0_m1_m2_n_tblk_lens = xdlops_gemm.GetCM0M1M2NThreadBlkLengths();
 
@@ -136,34 +136,34 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
         return make_naive_tensor_descriptor_packed(make_tuple(I1, I1, I1, I1, M0, M1, M2, N));
     }
 
-    __host__ __device__ static constexpr auto GetCM0N0M1N1M2M3M4N2BlockDescriptor()
+    __host__ __device__ static constexpr auto GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2()
     {
-        constexpr auto c_m0_n0_m1_n1_m2_n2_block_desc =
+        constexpr auto c_block_desc_m0_n0_m1_n1_m2_n2 =
             make_naive_tensor_descriptor_packed(make_tuple(Number<MRepeat>{},
                                                            Number<NRepeat>{},
                                                            Number<MWaves>{},
                                                            Number<NWaves>{},
                                                            Number<MPerXDL>{},
                                                            Number<NPerXDL>{}));
 
-        return xdlops_gemm.MakeCM0N0M1N1M2M3M4N2Descriptor(c_m0_n0_m1_n1_m2_n2_block_desc);
+        return xdlops_gemm.MakeCDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(c_block_desc_m0_n0_m1_n1_m2_n2);
     }
 
-    template <typename CMNGridDesc>
+    template <typename CGridDesc_M_N>
     __host__ __device__ static constexpr auto
-    MakeCM0N0M1N1M2M3M4N2GridDescriptor(const CMNGridDesc& c_m_n_grid_desc)
+    MakeCGridDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(const CGridDesc_M_N& c_grid_desc_m_n)
     {
-        const auto c_m0_n0_m1_n1_m2_n2_grid_desc = transform_tensor_descriptor(
-            c_m_n_grid_desc,
+        const auto c_grid_desc_m0_n0_m1_n1_m2_n2 = transform_tensor_descriptor(
+            c_grid_desc_m_n,
             make_tuple(make_unmerge_transform(make_tuple(MRepeat, MWaves, MPerXDL)),
                        make_unmerge_transform(make_tuple(NRepeat, NWaves, NPerXDL))),
             make_tuple(Sequence<0>{}, Sequence<1>{}),
             make_tuple(Sequence<0, 2, 4>{}, Sequence<1, 3, 5>{}));
 
-        return xdlops_gemm.MakeCM0N0M1N1M2M3M4N2Descriptor(c_m0_n0_m1_n1_m2_n2_grid_desc);
+        return xdlops_gemm.MakeCDescriptor_M0_N0_M1_N1_M2_M3_M4_N2(c_grid_desc_m0_n0_m1_n1_m2_n2);
     }
 
-    __host__ __device__ static constexpr auto MakeAK0M0M1M2K1BlockDescriptor()
+    __host__ __device__ static constexpr auto MakeABlockDescriptor_K0_M0_M1_M2_K1()
     {
         return transform_tensor_descriptor(
             AK0MK1BlockDesc{},
@@ -175,7 +175,7 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
             make_tuple(Sequence<0>{}, Sequence<1, 2, 3>{}, Sequence<4>{}));
     }
 
-    __host__ __device__ static constexpr auto MakeBK0N0N1N2K1BlockDescriptor()
+    __host__ __device__ static constexpr auto MakeBBlockDescriptor_K0_N0_N1_N2_K1()
     {
         return transform_tensor_descriptor(
             BK0NK1BlockDesc{},
@@ -187,8 +187,8 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
             make_tuple(Sequence<0>{}, Sequence<1, 2, 3>{}, Sequence<4>{}));
     }
 
-    static constexpr auto a_k0_m0_m1_m2_k1_block_desc = MakeAK0M0M1M2K1BlockDescriptor();
-    static constexpr auto b_k0_n0_n1_n2_k1_block_desc = MakeBK0N0N1N2K1BlockDescriptor();
+    static constexpr auto a_block_desc_k0_m0_m1_m2_k1 = MakeABlockDescriptor_K0_M0_M1_M2_K1();
+    static constexpr auto b_block_desc_k0_n0_n1_n2_k1 = MakeBBlockDescriptor_K0_N0_N1_N2_K1();
 
     template <typename ABlockBuffer, typename BBlockBuffer, typename CThreadBuffer>
     __device__ void Run(const ABlockBuffer& a_block_buf,
@@ -202,7 +202,7 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
 
         static_for<0, MRepeat, 1>{}([&](auto m0) {
             // read A
-            a_thread_copy_.Run(a_k0_m0_m1_m2_k1_block_desc,
+            a_thread_copy_.Run(a_block_desc_k0_m0_m1_m2_k1,
                                make_tuple(I0, m0, I0, I0, I0),
                                a_block_buf,
                                a_thread_desc_,
@@ -211,7 +211,7 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
 
             static_for<0, NRepeat, 1>{}([&](auto n0) {
                 // read B
-                b_thread_copy_.Run(b_k0_n0_n1_n2_k1_block_desc,
+                b_thread_copy_.Run(b_block_desc_k0_n0_n1_n2_k1,
                                    make_tuple(I0, n0, I0, I0, I0),
                                    b_block_buf,
                                    b_thread_desc_,
@@ -256,7 +256,7 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
 
     using AThreadCopy = ThreadwiseTensorSliceTransfer_v4<FloatAB,
                                                          FloatAB,
-                                                         decltype(a_k0_m0_m1_m2_k1_block_desc),
+                                                         decltype(a_block_desc_k0_m0_m1_m2_k1),
                                                          decltype(a_thread_desc_),
                                                          Sequence<K0, 1, 1, 1, K1>,
                                                          Sequence<0, 1, 2, 3, 4>,
@@ -266,7 +266,7 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
 
     using BThreadCopy = ThreadwiseTensorSliceTransfer_v4<FloatAB,
                                                          FloatAB,
-                                                         decltype(b_k0_n0_n1_n2_k1_block_desc),
+                                                         decltype(b_block_desc_k0_n0_n1_n2_k1),
                                                          decltype(b_thread_desc_),
                                                          Sequence<K0, 1, 1, 1, K1>,
                                                          Sequence<0, 1, 2, 3, 4>,