Move device library to seperate directory

src/targets/gpu/CMakeLists.txt

@@ -11,11 +11,12 @@ if(NOT TARGET MIOpen)
 endif()
 
 add_library(migraph_device 
-    hip_contiguous.cpp
+    device/contiguous.cpp
 )
 rocm_clang_tidy_check(migraph_device)
 target_link_libraries(migraph_device migraph hip::device)
 target_include_directories(migraph_device PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
+target_include_directories(migraph_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/device>)
 
 add_library(migraph_gpu
     eliminate_allocation.cpp

src/targets/gpu/hip_contiguous.cpp → src/targets/gpu/device/contiguous.cpp

 
 #include <hip/hip_runtime.h>
-#include <migraph/operators.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
 
 namespace migraph {
 namespace gpu {
+namespace device {
 
 struct index
 {
@@ -107,7 +108,7 @@ struct hip_tensor_descriptor
     size_t strides[NDim] = {};
 };
 
-void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph::argument result)
+void contiguous(shape output_shape, argument arg, argument result)
 {
     visit_all(result, arg)([&](auto output, auto input) {
         visit_tensor_size(output_shape.lens().size(), [&](auto ndim) {
@@ -130,5 +131,6 @@ void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph:
         });
     });
 }
+} // namespace device
 } // namespace gpu
 } // namespace migraph

src/targets/gpu/include/migraph/gpu/kernels.hpp → src/targets/gpu/include/migraph/gpu/device/contiguous.hpp

@@ -5,11 +5,12 @@
 
 namespace migraph {
 namespace gpu {
+namespace device {
 
-void hip_contiguous(shape output_shape, argument arg, argument result);
+void contiguous(shape output_shape, argument arg, argument result);
 
+} // namespace device
 } // namespace gpu
-
 } // namespace migraph
 
 #endif

src/targets/gpu/lowering.cpp

@@ -8,7 +8,7 @@
 #include <migraph/gpu/miopen.hpp>
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
-#include <migraph/gpu/kernels.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
 #include <migraph/iterator_for.hpp>
 #include <migraph/gpu/rocblas.hpp>
 #include <migraph/gpu/context.hpp>
@@ -252,7 +252,7 @@ struct miopen_contiguous
     {
         assert(output_shape == args[1].get_shape());
         assert(output_shape.standard());
-        hip_contiguous(std::move(output_shape), args.at(0), args.at(1));
+        device::contiguous(std::move(output_shape), args.at(0), args.at(1));
         return args.at(1);
     }
 };