From 81497a93a0840d5a1b5e84c1e47a90ae39d0fee6 Mon Sep 17 00:00:00 2001
From: Chao Liu <lc.roy86@gmail.com>
Date: Tue, 11 Jun 2019 23:49:51 -0500
Subject: [PATCH] reorginze files

---
 CMakeLists.txt                                 |  5 ++++-
 ...ce_convolution_direct_v2_nchw_kcyx_nkhw.hpp |  2 +-
 ...olution_implicit_gemm_v1_chwn_cyxk_khwn.hpp |  8 ++++----
 ...olution_implicit_gemm_v1_nchw_cyxk_nkhw.hpp |  4 ++--
 ...olution_implicit_gemm_v2_chwn_cyxk_khwn.hpp |  4 ++--
 ...olution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp |  4 ++--
 ...olution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp |  4 ++--
 ...convolution_2_vectorized_nchw_kcyx_nkhw.hpp |  2 +-
 ...emm_convolution_1_chwn_cyxk_khwn_padded.hpp |  2 +-
 driver/driver.cpp                              |  4 ++--
 ...se_convolution_direct_v2_nchw_kcyx_nkhw.hpp | 12 ++++++------
 ...ution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp | 16 ++++++++--------
 ...ution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp | 18 +++++++++---------
 ...ution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp | 16 ++++++++--------
 ...m_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp | 16 ++++++++--------
 ...ution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp | 16 ++++++++--------
 ...m_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp | 16 ++++++++--------
 ...olution_implicit_gemm_v2_chwn_cyxk_khwn.hpp | 12 ++++++------
 ...emm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp | 14 +++++++-------
 ...olution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp | 12 ++++++------
 ...emm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp | 12 ++++++------
 ...olution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp | 14 +++++++-------
 ...emm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp | 14 +++++++-------
 ...convolution_2_vectorized_nchw_kcyx_nkhw.hpp | 14 +++++++-------
 ...emm_convolution_1_chwn_cyxk_khwn_padded.hpp | 14 +++++++-------
 .../ConstantMatrixDescriptor.hpp               |  2 +-
 .../ConstantMergedTensorDescriptor.hpp         |  4 ++--
 .../ConstantTensorDescriptor.hpp               |  2 +-
 .../blockwise_2d_tensor_op.hpp                 |  4 ++--
 .../blockwise_3d_tensor_op.hpp                 |  4 ++--
 .../blockwise_4d_tensor_op.hpp                 |  4 ++--
 .../blockwise_batched_gemm.hpp                 |  2 +-
 .../tensor_operation}/blockwise_gemm.hpp       |  4 ++--
 .../blockwise_generic_tensor_slice_copy.hpp    |  2 +-
 .../blockwise_tensor_slice_copy.hpp            |  2 +-
 .../threadwise_4d_tensor_op.hpp                |  2 +-
 .../threadwise_direct_convolution.hpp          |  4 ++--
 .../tensor_operation}/threadwise_gemm.hpp      |  4 ++--
 .../threadwise_generic_tensor_op.hpp           |  4 ++--
 .../threadwise_generic_tensor_slice_copy.hpp   |  4 ++--
 .../threadwise_tensor_slice_copy.hpp           |  2 +-
 .../composable_kernel/utility}/Array.hpp       |  4 ++--
 .../composable_kernel/utility}/Sequence.hpp    |  4 ++--
 .../utility}/amd_inline_asm.hpp                |  2 +-
 include/composable_kernel/utility/common.hpp   | 17 +++++++++++++++++
 .../composable_kernel/utility}/config.hpp.in   |  0
 .../composable_kernel/utility}/functional.hpp  |  8 ++++----
 .../composable_kernel/utility}/functional2.hpp |  4 ++--
 .../composable_kernel/utility}/functional3.hpp |  8 ++++----
 .../utility}/integral_constant.hpp             |  0
 .../composable_kernel/utility}/utility.hpp     |  9 +++++----
 .../composable_kernel/utility}/vector_type.hpp |  4 ++--
 {src/include => include}/conv_common.hpp       |  2 +-
 {src/include => include}/device.hpp            |  2 +-
 .../gridwise_convolution_kernel_wrapper.hpp    |  4 ----
 {src/include => include}/tensor.hpp            |  0
 src/CMakeLists.txt                             |  3 +--
 src/device.cpp                                 |  2 +-
 src/include/common.hpp                         | 17 -----------------
 59 files changed, 197 insertions(+), 198 deletions(-)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp (96%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp (96%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp (96%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp (96%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp (97%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp (95%)
 rename {src/include => include/composable_kernel/kernel_algorithm}/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp (96%)
 rename {src/include => include/composable_kernel/tensor_description}/ConstantMatrixDescriptor.hpp (97%)
 rename {src/include => include/composable_kernel/tensor_description}/ConstantMergedTensorDescriptor.hpp (98%)
 rename {src/include => include/composable_kernel/tensor_description}/ConstantTensorDescriptor.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_2d_tensor_op.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_3d_tensor_op.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_4d_tensor_op.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_batched_gemm.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_gemm.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_generic_tensor_slice_copy.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/blockwise_tensor_slice_copy.hpp (99%)
 rename {src/include => include/composable_kernel/tensor_operation}/threadwise_4d_tensor_op.hpp (95%)
 rename {src/include => include/composable_kernel/tensor_operation}/threadwise_direct_convolution.hpp (98%)
 rename {src/include => include/composable_kernel/tensor_operation}/threadwise_gemm.hpp (96%)
 rename {src/include => include/composable_kernel/tensor_operation}/threadwise_generic_tensor_op.hpp (73%)
 rename {src/include => include/composable_kernel/tensor_operation}/threadwise_generic_tensor_slice_copy.hpp (96%)
 rename {src/include => include/composable_kernel/tensor_operation}/threadwise_tensor_slice_copy.hpp (99%)
 rename {src/include => include/composable_kernel/utility}/Array.hpp (99%)
 rename {src/include => include/composable_kernel/utility}/Sequence.hpp (99%)
 rename {src/include => include/composable_kernel/utility}/amd_inline_asm.hpp (99%)
 create mode 100644 include/composable_kernel/utility/common.hpp
 rename {src/include => include/composable_kernel/utility}/config.hpp.in (100%)
 rename {src/include => include/composable_kernel/utility}/functional.hpp (90%)
 rename {src/include => include/composable_kernel/utility}/functional2.hpp (93%)
 rename {src/include => include/composable_kernel/utility}/functional3.hpp (93%)
 rename {src/include => include/composable_kernel/utility}/integral_constant.hpp (100%)
 rename {src/include => include/composable_kernel/utility}/utility.hpp (93%)
 rename {src/include => include/composable_kernel/utility}/vector_type.hpp (96%)
 rename {src/include => include}/conv_common.hpp (98%)
 rename {src/include => include}/device.hpp (96%)
 rename {src/include => include}/gridwise_convolution_kernel_wrapper.hpp (93%)
 rename {src/include => include}/tensor.hpp (100%)
 delete mode 100644 src/include/common.hpp

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 252807ff4c..fbbafb877d 100644
--- a/CMakeLists.txt
+++ b/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)
diff --git a/driver/device_convolution_direct_v2_nchw_kcyx_nkhw.hpp b/driver/device_convolution_direct_v2_nchw_kcyx_nkhw.hpp
index 8e1126bd77..de1353a4ed 100644
--- a/driver/device_convolution_direct_v2_nchw_kcyx_nkhw.hpp
+++ b/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;
 
diff --git a/driver/device_convolution_implicit_gemm_v1_chwn_cyxk_khwn.hpp b/driver/device_convolution_implicit_gemm_v1_chwn_cyxk_khwn.hpp
index 4ffa1de4d9..67d27ec273 100644
--- a/driver/device_convolution_implicit_gemm_v1_chwn_cyxk_khwn.hpp
+++ b/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;
 
diff --git a/driver/device_convolution_implicit_gemm_v1_nchw_cyxk_nkhw.hpp b/driver/device_convolution_implicit_gemm_v1_nchw_cyxk_nkhw.hpp
index b973d19237..5c2dff2fb4 100644
--- a/driver/device_convolution_implicit_gemm_v1_nchw_cyxk_nkhw.hpp
+++ b/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;
 
diff --git a/driver/device_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp b/driver/device_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
index 09c29fce7b..7575d061c1 100644
--- a/driver/device_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
+++ b/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;
 
diff --git a/driver/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp b/driver/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
index 3a7ea0ab7b..2ac490e7fa 100644
--- a/driver/device_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
+++ b/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;
 
diff --git a/driver/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp b/driver/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
index 44434d46e8..92213a3029 100644
--- a/driver/device_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
+++ b/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;
 
diff --git a/driver/device_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp b/driver/device_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
index 386abe5ddd..0911607032 100644
--- a/driver/device_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
+++ b/driver/device_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
@@ -1,7 +1,7 @@
 #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;
 
diff --git a/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp b/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
index 68011d8d9a..98d55e1538 100644
--- a/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
+++ b/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
@@ -1,7 +1,7 @@
 #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;
 
diff --git a/driver/driver.cpp b/driver/driver.cpp
index 636a1b1109..4c6d9474ff 100644
--- a/driver/driver.cpp
+++ b/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"
diff --git a/src/include/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp
similarity index 96%
rename from src/include/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp
index 83b2a67f7f..f788bdd093 100644
--- a/src/include/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_direct_v2_nchw_kcyx_nkhw.hpp
@@ -1,12 +1,12 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp
index 8c42441fac..711f3464b6 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r1_chwn_cyxk_khwn.hpp
@@ -1,14 +1,14 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp
similarity index 96%
rename from src/include/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp
index 6d7bc47d37..6a3471c7f6 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r2_chwn_cyxk_khwn.hpp
@@ -1,15 +1,15 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp
index 5f5d8a1d92..5f6ba76d4b 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn.hpp
@@ -1,14 +1,14 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp
index 3ef597cb56..3f96d2de9d 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_chwn_cyxk_khwn_lds_double_buffer.hpp
@@ -1,14 +1,14 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp
index 4b13903459..a685727b50 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw.hpp
@@ -1,14 +1,14 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp
index 885252cd38..d531f5c0b2 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v1r3_nchw_cyxk_nkhw_lds_double_buffer.hpp
@@ -1,14 +1,14 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
similarity index 96%
rename from src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
index 26c6b75fa9..dd7e02e715 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn.hpp
@@ -1,12 +1,12 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp
index 3241aabf62..9839f7de24 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v2_chwn_cyxk_khwn_lds_double_buffer.hpp
@@ -1,13 +1,13 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
index ab44486280..67440af681 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw.hpp
@@ -1,12 +1,12 @@
 #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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp
index 2222e53abf..d3cdb640ff 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v3_nchw_cyxk_nkhw_lds_double_buffer.hpp
@@ -1,12 +1,12 @@
 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V3_NCHW_CYXK_NKHW_LDS_DOUBLE_BUFFER
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V3_NCHW_CYXK_NKHW_LDS_DOUBLE_BUFFER
 
-#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 {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
similarity index 96%
rename from src/include/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
index 31067c8591..c97133dfad 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw.hpp
@@ -1,13 +1,13 @@
 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V4_NCHW_KCYX_NKHW
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V4_NCHW_KCYX_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 "threadwise_generic_tensor_slice_copy.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"
+#include "composable_kernel/tensor_operation/threadwise_generic_tensor_slice_copy.hpp"
 
 namespace ck {
 
diff --git a/src/include/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp
similarity index 97%
rename from src/include/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp
index 9df6700d6d..786d83db97 100644
--- a/src/include/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_convolution_implicit_gemm_v4_nchw_kcyx_nkhw_lds_double_buffer.hpp
@@ -1,13 +1,13 @@
 #ifndef CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V4_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER
 #define CK_GRIDWISE_CONVOLUTION_IMPLICIT_GEMM_V4_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMergedTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_generic_tensor_slice_copy.hpp"
-#include "blockwise_gemm.hpp"
-#include "threadwise_generic_tensor_slice_copy.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"
+#include "composable_kernel/tensor_operation/threadwise_generic_tensor_slice_copy.hpp"
 
 namespace ck {
 
diff --git a/src/include/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp b/include/composable_kernel/kernel_algorithm/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
similarity index 95%
rename from src/include/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
index f2552260bc..e160ef8686 100644
--- a/src/include/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_direct_convolution_2_vectorized_nchw_kcyx_nkhw.hpp
@@ -1,11 +1,11 @@
 #pragma once
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "blockwise_direct_convolution.hpp"
-#include "threadwise_4d_tensor_op.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/blockwise_direct_convolution.hpp"
+#include "composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/threadwise_direct_convolution.hpp"
 
 namespace ck {
 
diff --git a/src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp b/include/composable_kernel/kernel_algorithm/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
similarity index 96%
rename from src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
rename to include/composable_kernel/kernel_algorithm/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
index 0dd0dc1f20..c4dfb03dac 100644
--- a/src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
+++ b/include/composable_kernel/kernel_algorithm/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp
@@ -1,11 +1,11 @@
 #pragma once
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMatrixDescriptor.hpp"
-#include "blockwise_4d_tensor_op.hpp"
-#include "blockwise_2d_tensor_op.hpp"
-#include "threadwise_4d_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/threadwise_4d_tensor_op.hpp"
+#include "composable_kernel/tensor_operation/blockwise_gemm.hpp"
 
 namespace ck {
 
diff --git a/src/include/ConstantMatrixDescriptor.hpp b/include/composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp
similarity index 97%
rename from src/include/ConstantMatrixDescriptor.hpp
rename to include/composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp
index a2b88001ac..36b6efc467 100644
--- a/src/include/ConstantMatrixDescriptor.hpp
+++ b/include/composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_CONSTANT_MATRIX_DESCRIPTOR_HPP
 #define CK_CONSTANT_MATRIX_DESCRIPTOR_HPP
 
-#include "common.hpp"
+#include "composable_kernel/utility/common.hpp"
 
 namespace ck {
 
diff --git a/src/include/ConstantMergedTensorDescriptor.hpp b/include/composable_kernel/tensor_description/ConstantMergedTensorDescriptor.hpp
similarity index 98%
rename from src/include/ConstantMergedTensorDescriptor.hpp
rename to include/composable_kernel/tensor_description/ConstantMergedTensorDescriptor.hpp
index 244bd0eec7..600577b07b 100644
--- a/src/include/ConstantMergedTensorDescriptor.hpp
+++ b/include/composable_kernel/tensor_description/ConstantMergedTensorDescriptor.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_CONSTANT_MERGED_TENSOR_DESCRIPTOR_HPP
 #define CK_CONSTANT_MERGED_TENSOR_DESCRIPTOR_HPP
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/ConstantTensorDescriptor.hpp b/include/composable_kernel/tensor_description/ConstantTensorDescriptor.hpp
similarity index 99%
rename from src/include/ConstantTensorDescriptor.hpp
rename to include/composable_kernel/tensor_description/ConstantTensorDescriptor.hpp
index 5325259858..c9415cd9d4 100644
--- a/src/include/ConstantTensorDescriptor.hpp
+++ b/include/composable_kernel/tensor_description/ConstantTensorDescriptor.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_CONSTANT_TENSOR_DESCRIPTOR_HPP
 #define CK_CONSTANT_TENSOR_DESCRIPTOR_HPP
 
-#include "common.hpp"
+#include "composable_kernel/utility/common.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_2d_tensor_op.hpp b/include/composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp
similarity index 99%
rename from src/include/blockwise_2d_tensor_op.hpp
rename to include/composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp
index aeb4738d7e..09ef7476cf 100644
--- a/src/include/blockwise_2d_tensor_op.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_2d_tensor_op.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_BLOCKWISE_2D_TENSOR_OP_HPP
 #define CK_BLOCKWISE_2D_TENSOR_OP_HPP
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_3d_tensor_op.hpp b/include/composable_kernel/tensor_operation/blockwise_3d_tensor_op.hpp
similarity index 99%
rename from src/include/blockwise_3d_tensor_op.hpp
rename to include/composable_kernel/tensor_operation/blockwise_3d_tensor_op.hpp
index 4633dc1075..e17972a1a9 100644
--- a/src/include/blockwise_3d_tensor_op.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_3d_tensor_op.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_BLOCKWISE_3D_TENSOR_OP_HPP
 #define CK_BLOCKWISE_3D_TENSOR_OP_HPP
 
-#include "common.hpp"
-#include "ConstantTensorDescriptor.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_4d_tensor_op.hpp b/include/composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp
similarity index 99%
rename from src/include/blockwise_4d_tensor_op.hpp
rename to include/composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp
index 51c30b08e4..bad94e6f35 100644
--- a/src/include/blockwise_4d_tensor_op.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_4d_tensor_op.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_BLOCKWISE_4D_TENSOR_OP_HPP
 #define CK_BLOCKWISE_4D_TENSOR_OP_HPP
 
-#include "ConstantTensorDescriptor.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_batched_gemm.hpp b/include/composable_kernel/tensor_operation/blockwise_batched_gemm.hpp
similarity index 99%
rename from src/include/blockwise_batched_gemm.hpp
rename to include/composable_kernel/tensor_operation/blockwise_batched_gemm.hpp
index 340c25df55..db51bb0b1d 100644
--- a/src/include/blockwise_batched_gemm.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_batched_gemm.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_BLOCKWISE_BATCHED_GEMM_HPP
 #define CK_BLOCKWISE_BATCHED_GEMM_HPP
 
-#include "threadwise_gemm.hpp"
+#include "composable_kernel/tensor_operation/threadwise_gemm.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_gemm.hpp b/include/composable_kernel/tensor_operation/blockwise_gemm.hpp
similarity index 99%
rename from src/include/blockwise_gemm.hpp
rename to include/composable_kernel/tensor_operation/blockwise_gemm.hpp
index 0fc9a7bb67..c045c32c03 100644
--- a/src/include/blockwise_gemm.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_gemm.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_BLOCKWISE_GEMM_HPP
 #define CK_BLOCKWISE_GEMM_HPP
 
-#include "common.hpp"
-#include "threadwise_gemm.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_operation/threadwise_gemm.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_generic_tensor_slice_copy.hpp b/include/composable_kernel/tensor_operation/blockwise_generic_tensor_slice_copy.hpp
similarity index 99%
rename from src/include/blockwise_generic_tensor_slice_copy.hpp
rename to include/composable_kernel/tensor_operation/blockwise_generic_tensor_slice_copy.hpp
index 6b887380e7..1381b7ef14 100644
--- a/src/include/blockwise_generic_tensor_slice_copy.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_generic_tensor_slice_copy.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_BLOCKWISE_GENERIC_TENSOR_SLICE_COPY_HPP
 #define CK_BLOCKWISE_GENERIC_TENSOR_SLICE_COPY_HPP
 
-#include "threadwise_generic_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_generic_tensor_slice_copy.hpp"
 
 namespace ck {
 
diff --git a/src/include/blockwise_tensor_slice_copy.hpp b/include/composable_kernel/tensor_operation/blockwise_tensor_slice_copy.hpp
similarity index 99%
rename from src/include/blockwise_tensor_slice_copy.hpp
rename to include/composable_kernel/tensor_operation/blockwise_tensor_slice_copy.hpp
index ed0f12c9a5..d4576279a8 100644
--- a/src/include/blockwise_tensor_slice_copy.hpp
+++ b/include/composable_kernel/tensor_operation/blockwise_tensor_slice_copy.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_BLOCKWISE_TENSOR_SLICE_COPY_HPP
 #define CK_BLOCKWISE_TENSOR_SLICE_COPY_HPP
 
-#include "threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
 
 namespace ck {
 
diff --git a/src/include/threadwise_4d_tensor_op.hpp b/include/composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp
similarity index 95%
rename from src/include/threadwise_4d_tensor_op.hpp
rename to include/composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp
index 9f8f0d12d6..1af6e4dc34 100644
--- a/src/include/threadwise_4d_tensor_op.hpp
+++ b/include/composable_kernel/tensor_operation/threadwise_4d_tensor_op.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_THREADWISE_4D_TENSOR_OP_HPP
 #define CK_THREADWISE_4D_TENSOR_OP_HPP
 
-#include "ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/threadwise_direct_convolution.hpp b/include/composable_kernel/tensor_operation/threadwise_direct_convolution.hpp
similarity index 98%
rename from src/include/threadwise_direct_convolution.hpp
rename to include/composable_kernel/tensor_operation/threadwise_direct_convolution.hpp
index 5d4d6376bb..0332dd9d3d 100644
--- a/src/include/threadwise_direct_convolution.hpp
+++ b/include/composable_kernel/tensor_operation/threadwise_direct_convolution.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_THREADWISE_DIRECT_CONVOLUTION_HPP
 #define CK_THREADWISE_DIRECT_CONVOLUTION_HPP
 
-#include "ConstantTensorDescriptor.hpp"
-#include "threadwise_tensor_slice_copy.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp"
 
 namespace ck {
 
diff --git a/src/include/threadwise_gemm.hpp b/include/composable_kernel/tensor_operation/threadwise_gemm.hpp
similarity index 96%
rename from src/include/threadwise_gemm.hpp
rename to include/composable_kernel/tensor_operation/threadwise_gemm.hpp
index d79e2e9f9c..5610b65e09 100644
--- a/src/include/threadwise_gemm.hpp
+++ b/include/composable_kernel/tensor_operation/threadwise_gemm.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_THREADWISE_GEMM_HPP
 #define CK_THREADWISE_GEMM_HPP
 
-#include "common.hpp"
-#include "ConstantMatrixDescriptor.hpp"
+#include "composable_kernel/utility/common.hpp"
+#include "composable_kernel/tensor_description/ConstantMatrixDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/threadwise_generic_tensor_op.hpp b/include/composable_kernel/tensor_operation/threadwise_generic_tensor_op.hpp
similarity index 73%
rename from src/include/threadwise_generic_tensor_op.hpp
rename to include/composable_kernel/tensor_operation/threadwise_generic_tensor_op.hpp
index e402255865..beca415604 100644
--- a/src/include/threadwise_generic_tensor_op.hpp
+++ b/include/composable_kernel/tensor_operation/threadwise_generic_tensor_op.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_THREADWISE_GENERIC_TENSOR_OP_HPP
 #define CK_THREADWISE_GENERIC_TENSOR_OP_HPP
 
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMergedTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMergedTensorDescriptor.hpp"
 
 namespace ck {
 template <class Float, class TDesc>
diff --git a/src/include/threadwise_generic_tensor_slice_copy.hpp b/include/composable_kernel/tensor_operation/threadwise_generic_tensor_slice_copy.hpp
similarity index 96%
rename from src/include/threadwise_generic_tensor_slice_copy.hpp
rename to include/composable_kernel/tensor_operation/threadwise_generic_tensor_slice_copy.hpp
index 9e7f90fe60..588766b477 100644
--- a/src/include/threadwise_generic_tensor_slice_copy.hpp
+++ b/include/composable_kernel/tensor_operation/threadwise_generic_tensor_slice_copy.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_THREADWISE_GENERIC_TENSOR_SLICE_COPY_HPP
 #define CK_THREADWISE_GENERIC_TENSOR_SLICE_COPY_HPP
 
-#include "ConstantTensorDescriptor.hpp"
-#include "ConstantMergedTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantMergedTensorDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/threadwise_tensor_slice_copy.hpp b/include/composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp
similarity index 99%
rename from src/include/threadwise_tensor_slice_copy.hpp
rename to include/composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp
index a32e7e3936..d48eed6ad1 100644
--- a/src/include/threadwise_tensor_slice_copy.hpp
+++ b/include/composable_kernel/tensor_operation/threadwise_tensor_slice_copy.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_THREADWISE_TENSOR_SLICE_COPY_HPP
 #define CK_THREADWISE_TENSOR_SLICE_COPY_HPP
 
-#include "ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 
 namespace ck {
 
diff --git a/src/include/Array.hpp b/include/composable_kernel/utility/Array.hpp
similarity index 99%
rename from src/include/Array.hpp
rename to include/composable_kernel/utility/Array.hpp
index fcf87c5843..f67f5a6cde 100644
--- a/src/include/Array.hpp
+++ b/include/composable_kernel/utility/Array.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_ARRAY_HPP
 #define CK_ARRAY_HPP
 
-#include "Sequence.hpp"
-#include "functional2.hpp"
+#include "composable_kernel/utility/Sequence.hpp"
+#include "composable_kernel/utility/functional2.hpp"
 
 namespace ck {
 
diff --git a/src/include/Sequence.hpp b/include/composable_kernel/utility/Sequence.hpp
similarity index 99%
rename from src/include/Sequence.hpp
rename to include/composable_kernel/utility/Sequence.hpp
index 09d67c13fa..70c3a5dc39 100644
--- a/src/include/Sequence.hpp
+++ b/include/composable_kernel/utility/Sequence.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_SEQUENCE_HPP
 #define CK_SEQUENCE_HPP
 
-#include "integral_constant.hpp"
-#include "functional.hpp"
+#include "composable_kernel/utility/integral_constant.hpp"
+#include "composable_kernel/utility/functional.hpp"
 
 namespace ck {
 
diff --git a/src/include/amd_inline_asm.hpp b/include/composable_kernel/utility/amd_inline_asm.hpp
similarity index 99%
rename from src/include/amd_inline_asm.hpp
rename to include/composable_kernel/utility/amd_inline_asm.hpp
index 70e8cf31ba..344bc87063 100644
--- a/src/include/amd_inline_asm.hpp
+++ b/include/composable_kernel/utility/amd_inline_asm.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_AMD_INLINE_ASM_HPP
 #define CK_AMD_INLINE_ASM_HPP
 
-#include "common.hpp"
+#include "composable_kernel/utility/vector_type.hpp"
 
 #define NO_VM_WAIT 0
 #define NO_LGKM_WAIT 0
diff --git a/include/composable_kernel/utility/common.hpp b/include/composable_kernel/utility/common.hpp
new file mode 100644
index 0000000000..08167cbccc
--- /dev/null
+++ b/include/composable_kernel/utility/common.hpp
@@ -0,0 +1,17 @@
+#ifndef CK_COMMON_HPP
+#define CK_COMMON_HPP
+
+#include "composable_kernel/utility/utility.hpp"
+#include "composable_kernel/utility/vector_type.hpp"
+#include "composable_kernel/utility/integral_constant.hpp"
+#include "composable_kernel/utility/Sequence.hpp"
+#include "composable_kernel/utility/Array.hpp"
+#include "composable_kernel/utility/functional.hpp"
+#include "composable_kernel/utility/functional2.hpp"
+#include "composable_kernel/utility/functional3.hpp"
+
+#if CK_USE_AMD_INLINE_ASM
+#include "composable_kernel/utility/amd_inline_asm.hpp"
+#endif
+
+#endif
diff --git a/src/include/config.hpp.in b/include/composable_kernel/utility/config.hpp.in
similarity index 100%
rename from src/include/config.hpp.in
rename to include/composable_kernel/utility/config.hpp.in
diff --git a/src/include/functional.hpp b/include/composable_kernel/utility/functional.hpp
similarity index 90%
rename from src/include/functional.hpp
rename to include/composable_kernel/utility/functional.hpp
index 15c957e1c5..92102ca418 100644
--- a/src/include/functional.hpp
+++ b/include/composable_kernel/utility/functional.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_FUNCTIONAL_HPP
 #define CK_FUNCTIONAL_HPP
 
-#include "integral_constant.hpp"
-#include "Sequence.hpp"
+#include "composable_kernel/utility/integral_constant.hpp"
+#include "composable_kernel/utility/Sequence.hpp"
 
 namespace ck {
 
@@ -38,7 +38,7 @@ struct static_if<true>
     __host__ __device__ constexpr auto operator()(F f) const
     {
         // This is a trick for compiler:
-        //   Pass forwarder to lambda "f" as "auto" argument, and maks sure "f" will use it,
+        //   Pass forwarder to lambda "f" as "auto" argument, and make sure "f" will use it,
         //   this will make "f" a generic lambda, so that "f" won't be compiled until being
         //   instantiated here
         f(forwarder{});
@@ -67,7 +67,7 @@ struct static_if<false>
     __host__ __device__ static constexpr auto Else(F f)
     {
         // This is a trick for compiler:
-        //   Pass forwarder to lambda "f" as "auto" argument, and maks sure "f" will use it,
+        //   Pass forwarder to lambda "f" as "auto" argument, and make sure "f" will use it,
         //   this will make "f" a generic lambda, so that "f" won't be compiled until being
         //   instantiated here
         f(forwarder{});
diff --git a/src/include/functional2.hpp b/include/composable_kernel/utility/functional2.hpp
similarity index 93%
rename from src/include/functional2.hpp
rename to include/composable_kernel/utility/functional2.hpp
index 3820056593..6d5e38d2ea 100644
--- a/src/include/functional2.hpp
+++ b/include/composable_kernel/utility/functional2.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_FUNCTIONAL2_HPP
 #define CK_FUNCTIONAL2_HPP
 
-#include "functional.hpp"
-#include "Sequence.hpp"
+#include "composable_kernel/utility/functional.hpp"
+#include "composable_kernel/utility/Sequence.hpp"
 
 namespace ck {
 
diff --git a/src/include/functional3.hpp b/include/composable_kernel/utility/functional3.hpp
similarity index 93%
rename from src/include/functional3.hpp
rename to include/composable_kernel/utility/functional3.hpp
index fc5f8a6bab..769f1bea0b 100644
--- a/src/include/functional3.hpp
+++ b/include/composable_kernel/utility/functional3.hpp
@@ -1,10 +1,10 @@
 #ifndef CK_FUNCTIONAL3_HPP
 #define CK_FUNCTIONAL3_HPP
 
-#include "functional.hpp"
-#include "functional2.hpp"
-#include "Sequence.hpp"
-#include "Array.hpp"
+#include "composable_kernel/utility/functional.hpp"
+#include "composable_kernel/utility/functional2.hpp"
+#include "composable_kernel/utility/Sequence.hpp"
+#include "composable_kernel/utility/Array.hpp"
 
 namespace ck {
 
diff --git a/src/include/integral_constant.hpp b/include/composable_kernel/utility/integral_constant.hpp
similarity index 100%
rename from src/include/integral_constant.hpp
rename to include/composable_kernel/utility/integral_constant.hpp
diff --git a/src/include/utility.hpp b/include/composable_kernel/utility/utility.hpp
similarity index 93%
rename from src/include/utility.hpp
rename to include/composable_kernel/utility/utility.hpp
index c7b4ba337d..ed834b0df0 100644
--- a/src/include/utility.hpp
+++ b/include/composable_kernel/utility/utility.hpp
@@ -1,5 +1,5 @@
-#ifndef CK_BASE_HPP
-#define CK_BASE_HPP
+#ifndef CK_UTILITY_HPP
+#define CK_UTILITY_HPP
 
 namespace ck {
 
@@ -25,7 +25,8 @@ __host__ __device__ constexpr bool is_same_type(X, Y)
     return is_same<X, Y>::value;
 }
 
-namespace math { // namespace math
+namespace math {
+
 template <class T, T s>
 struct scales
 {
@@ -106,7 +107,7 @@ __host__ __device__ constexpr T min(T x, Ts... xs)
 }
 
 // this is wrong
-// TODO: implement correct least common multiple, instead of calling max()
+// TODO: implement least common multiple properly, instead of calling max()
 template <class T, class... Ts>
 __host__ __device__ constexpr T lcm(T x, Ts... xs)
 {
diff --git a/src/include/vector_type.hpp b/include/composable_kernel/utility/vector_type.hpp
similarity index 96%
rename from src/include/vector_type.hpp
rename to include/composable_kernel/utility/vector_type.hpp
index 72c73068e1..7c48bdb1f8 100644
--- a/src/include/vector_type.hpp
+++ b/include/composable_kernel/utility/vector_type.hpp
@@ -1,8 +1,8 @@
 #ifndef CK_VECTOR_TYPE_HPP
 #define CK_VECTOR_TYPE_HPP
 
-#include "config.hpp"
-#include "integral_constant.hpp"
+#include "composable_kernel/utility/config.hpp"
+#include "composable_kernel/utility/integral_constant.hpp"
 
 namespace ck {
 
diff --git a/src/include/conv_common.hpp b/include/conv_common.hpp
similarity index 98%
rename from src/include/conv_common.hpp
rename to include/conv_common.hpp
index 254f4c5651..9f41a1d440 100644
--- a/src/include/conv_common.hpp
+++ b/include/conv_common.hpp
@@ -1,7 +1,7 @@
 #ifndef CK_CONV_COMMON_HPP
 #define CK_CONV_COMMON_HPP
 
-#include "ConstantTensorDescriptor.hpp"
+#include "composable_kernel/tensor_description/ConstantTensorDescriptor.hpp"
 
 using namespace ck;
 
diff --git a/src/include/device.hpp b/include/device.hpp
similarity index 96%
rename from src/include/device.hpp
rename to include/device.hpp
index 5766d8f990..2c982a4012 100644
--- a/src/include/device.hpp
+++ b/include/device.hpp
@@ -2,7 +2,7 @@
 #define CK_DEVICE_HPP
 
 #include <memory>
-#include "config.hpp"
+#include "composable_kernel/utility/config.hpp"
 
 using namespace ck;
 
diff --git a/src/include/gridwise_convolution_kernel_wrapper.hpp b/include/gridwise_convolution_kernel_wrapper.hpp
similarity index 93%
rename from src/include/gridwise_convolution_kernel_wrapper.hpp
rename to include/gridwise_convolution_kernel_wrapper.hpp
index a7caeed5aa..2269e72579 100644
--- a/src/include/gridwise_convolution_kernel_wrapper.hpp
+++ b/include/gridwise_convolution_kernel_wrapper.hpp
@@ -1,8 +1,6 @@
 #ifndef CK_GRIDWISE_CONVOLUTION_KERNEL_WRAPPER
 #define CK_GRIDWISE_CONVOLUTION_KERNEL_WRAPPER
 
-namespace ck {
-
 template <class GridwiseConvolution, class T>
 __global__ void run_gridwise_convolution_kernel(const T* const __restrict__ p_in_global,
                                                 const T* const __restrict__ p_wei_global,
@@ -11,6 +9,4 @@ __global__ void run_gridwise_convolution_kernel(const T* const __restrict__ p_in
     GridwiseConvolution{}.Run(p_in_global, p_wei_global, p_out_global);
 }
 
-} // namespace ck
-
 #endif
diff --git a/src/include/tensor.hpp b/include/tensor.hpp
similarity index 100%
rename from src/include/tensor.hpp
rename to include/tensor.hpp
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index a34b7b3c89..a0b63a179d 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -1,4 +1,4 @@
-configure_file("${PROJECT_SOURCE_DIR}/src/include/config.hpp.in" "${PROJECT_BINARY_DIR}/src/include/config.hpp")
+configure_file("${PROJECT_SOURCE_DIR}/include/composable_kernel/utility/config.hpp.in" "${PROJECT_BINARY_DIR}/include/composable_kernel/utility/config.hpp")
 
 set(TENSOR_SOURCE 
     tensor.cpp;
@@ -9,7 +9,6 @@ add_library(tensor SHARED ${TENSOR_SOURCE})
 target_compile_features(tensor PUBLIC)
 set_target_properties(tensor PROPERTIES POSITION_INDEPENDENT_CODE ON)
 
-
 if(DEVICE_BACKEND STREQUAL "CUDA")
     target_link_libraries(tensor nvToolsExt cudart)
 endif()
diff --git a/src/device.cpp b/src/device.cpp
index f65597ca76..ab880b33da 100644
--- a/src/device.cpp
+++ b/src/device.cpp
@@ -1,4 +1,4 @@
-#include "config.hpp"
+#include "composable_kernel/utility/config.hpp"
 #include "device.hpp"
 
 DeviceMem::DeviceMem(std::size_t mem_size) : mMemSize(mem_size)
diff --git a/src/include/common.hpp b/src/include/common.hpp
deleted file mode 100644
index e52aa7741f..0000000000
--- a/src/include/common.hpp
+++ /dev/null
@@ -1,17 +0,0 @@
-#ifndef CK_COMMON_HPP
-#define CK_COMMON_HPP
-
-#include "utility.hpp"
-#include "vector_type.hpp"
-#include "integral_constant.hpp"
-#include "Sequence.hpp"
-#include "Array.hpp"
-#include "functional.hpp"
-#include "functional2.hpp"
-#include "functional3.hpp"
-
-#if CK_USE_AMD_INLINE_ASM
-#include "amd_inline_asm.hpp"
-#endif
-
-#endif