reorginze files

CMakeLists.txt

@@ -45,6 +45,9 @@ elseif(DEVICE_BACKEND STREQUAL "CUDA")
 endif()
 
 #
-include_directories(BEFORE src/include ${PROJECT_BINARY_DIR}/src/include)
+include_directories(BEFORE
+    include
+    ${PROJECT_BINARY_DIR}/include
+)
 add_subdirectory(src)
 add_subdirectory(driver)

driver/device_convolution_direct_v2_nchw_kcyx_nkhw.hpp

@@ -2,7 +2,7 @@
 #include <unistd.h>
 #include "device.hpp"
 #include "gridwise_convolution_kernel_wrapper.hpp"
-#include "gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp"
 
 using namespace ck;
 

driver/device_convolution_implicit_gemm_v1_chwn_cyxk_khwn.hpp

@@ -2,10 +2,10 @@
 #include <unistd.h>
 #include "device.hpp"
 #include "gridwise_convolution_kernel_wrapper.hpp"
-#include "gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp"
-#include "gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp"
-#include "gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp"
-#include "gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp"
 
 using namespace ck;
 

driver/device_convolution_implicit_gemm_v1_nchw_cyxk_nkhw.hpp

@@ -2,8 +2,8 @@
 #include <unistd.h>
 #include "device.hpp"
 #include "gridwise_convolution_kernel_wrapper.hpp"
-#include "gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp"
-#include "gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp"
 
 using namespace ck;
 

driver/device_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp

@@ -2,8 +2,8 @@
 #include <unistd.h>
 #include "device.hpp"
 #include "gridwise_convolution_kernel_wrapper.hpp"
-#include "gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp"
-#include "gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp"
 
 using namespace ck;
 

driver/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp

@@ -2,8 +2,8 @@
 #include <unistd.h>
 #include "device.hpp"
 #include "gridwise_convolution_kernel_wrapper.hpp"
-#include "gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp"
-#include "gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp"
 
 using namespace ck;
 

driver/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp

@@ -2,8 +2,8 @@
 #include <unistd.h>
 #include "device.hpp"
 #include "gridwise_convolution_kernel_wrapper.hpp"
-#include "gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp"
-#include "gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp"
 
 using namespace ck;
 

driver/device_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp

 #pragma once
 #include <unistd.h>
 #include "device.hpp"
-#include "gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp"
 
 using namespace ck;
 

driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp

 #pragma once
 #include <unistd.h>
 #include "device.hpp"
-#include "gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp"
+#include "composable_kernel/kernel_algorithm/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp"
 
 using namespace ck;
 

driver/driver.cpp

@@ -3,9 +3,9 @@
 #include <initializer_list>
 #include <cstdlib>
 #include <stdlib.h>
-#include "config.hpp"
+#include "composable_kernel/utility/config.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 #include "tensor.hpp"
-#include "ConstantTensorDescriptor.hpp"
 #include "conv_common.hpp"
 #include "device_convolution_direct_v2_nchw_kcyx_nkhw.hpp"
 #include "device_convolution_implicit_gemm_v1_chwn_cyxk_khwn.hpp"

src/include/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_DIRECT_V2_NCHW_KCYX_NKHW
 #define CK_GRIDWISE_CONVOLUTION_DIRECT_V2_NCHW_KCYX_NKHW
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_direct_convolution.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_direct_convolution.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R1_CHWN_CYXK_KHWN
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R1_CHWN_CYXK_KHWN
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_4d_tensor_op.hpp"
-#include "blockwise_batched_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_batched_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R2_CHWN_CYXK_KHWN
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R2_CHWN_CYXK_KHWN
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_3d_tensor_op.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_4d_tensor_op.hpp"
-#include "blockwise_batched_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_3d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_batched_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_CHWN_CYXK_KHWN
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_CHWN_CYXK_KHWN
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_4d_tensor_op.hpp"
-#include "blockwise_batched_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_batched_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_CHWN_CYXK_KHWN_LDS_DOUBLE_BUFFER
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_CHWN_CYXK_KHWN_LDS_DOUBLE_BUFFER
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_4d_tensor_op.hpp"
-#include "blockwise_batched_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_batched_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_NCHW_CYXK_NKHW
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_NCHW_CYXK_NKHW
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_tensor_slice_copy.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_generic_tensor_op.hpp"
-#include "blockwise_batched_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_generic_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_batched_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_NCHW_CYXK_NKHW_LDS_DOUBLE_BUFFER
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V1R3_NCHW_CYXK_NKHW_LDS_DOUBLE_BUFFER
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_tensor_slice_copy.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "threadwise_generic_tensor_op.hpp"
-#include "blockwise_batched_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_generic_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_batched_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V2_CHWN_CYXK_KHWN
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V2_CHWN_CYXK_KHWN
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V2_CHWN_CYXK_KHWN_LDS_DOUBLE_BUFFER
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V2_CHWN_CYXK_KHWN_LDS_DOUBLE_BUFFER
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
-#include "blockwise_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/blockwise_gemm.hpp"
 
 namespace ck {
 

src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp → include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp

 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V3_NCHW_CYXK_NKHW
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V3_NCHW_CYXK_NKHW
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMergedTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_generic_tensor_slice_copy.hpp"
-#include "blockwise_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMergedTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/tensor_operation/blockwise_generic_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/blockwise_gemm.hpp"
 
 namespace ck {