[metal] Move platform/metal to backends/metal (taichi-dev#667)

archibate · Mar 31, 2020 · 9b79aab · 9b79aab
1 parent 7b2d372
commit 9b79aab
Show file tree

Hide file tree

Showing 22 changed files with 91 additions and 67 deletions.
diff --git a/cmake/TaichiCore.cmake b/cmake/TaichiCore.cmake
@@ -9,13 +9,17 @@ file(GLOB TAICHI_CORE_SOURCE
 file(GLOB TAICHI_BACKEND_SOURCE "taichi/backends/**/*.cpp" "taichi/backends/**/*.h")
 
 file(GLOB TAICHI_CUDA_SOURCE "taichi/backends/cuda/*.cpp" "taichi/backends/cuda/*.h")
+file(GLOB TAICHI_METAL_SOURCE "taichi/backends/metal/*.h" "taichi/backends/metal/*.cpp" "taichi/backends/metal/shaders/*")
 
 list(REMOVE_ITEM TAICHI_CORE_SOURCE ${TAICHI_BACKEND_SOURCE})
 
 if (TI_WITH_CUDA)
     list(APPEND TAICHI_CORE_SOURCE ${TAICHI_CUDA_SOURCE})
 endif()
 
+# TODO(#529) include Metal source only on Apple MacOS
+list(APPEND TAICHI_CORE_SOURCE ${TAICHI_METAL_SOURCE})
+
 option(BUILD_CPP_EXAMPLES "Build legacy C++ examples" OFF)
 
 if (BUILD_CPP_EXAMPLES)

diff --git a/taichi/platform/metal/metal_api.cpp → taichi/backends/metal/api.cpp b/taichi/platform/metal/metal_api.cpp → taichi/backends/metal/api.cpp
@@ -1,4 +1,4 @@
-#include "metal_api.h"
+#include "taichi/backends/metal/api.h"
 
 TLANG_NAMESPACE_BEGIN
 

diff --git a/taichi/platform/metal/metal_api.h → taichi/backends/metal/api.h b/taichi/platform/metal/metal_api.h → taichi/backends/metal/api.h
diff --git a/taichi/platform/metal/metal_data_types.cpp → taichi/backends/metal/data_types.cpp b/taichi/platform/metal/metal_data_types.cpp → taichi/backends/metal/data_types.cpp
@@ -1,4 +1,4 @@
-#include "metal_data_types.h"
+#include "taichi/backends/metal/data_types.h"
 
 TLANG_NAMESPACE_BEGIN
 

diff --git a/taichi/platform/metal/metal_data_types.h → taichi/backends/metal/data_types.h b/taichi/platform/metal/metal_data_types.h → taichi/backends/metal/data_types.h
@@ -1,8 +1,9 @@
 #pragma once
 
-#include <taichi/lang_util.h>
 #include <string>
 
+#include "taichi/lang_util.h"
+
 TLANG_NAMESPACE_BEGIN
 
 enum class MetalDataType : int {

diff --git a/taichi/platform/metal/metal_kernel_util.cpp → taichi/backends/metal/kernel_util.cpp b/taichi/platform/metal/metal_kernel_util.cpp → taichi/backends/metal/kernel_util.cpp
@@ -1,4 +1,4 @@
-#include "taichi/platform/metal/metal_kernel_util.h"
+#include "taichi/backends/metal/kernel_util.h"
 
 #define TI_RUNTIME_HOST
 #include "taichi/runtime/llvm/context.h"

diff --git a/taichi/platform/metal/metal_kernel_util.h → taichi/backends/metal/kernel_util.h b/taichi/platform/metal/metal_kernel_util.h → taichi/backends/metal/kernel_util.h
@@ -4,7 +4,7 @@
 #include <vector>
 
 #include "taichi/ir/statements.h"
-#include "taichi/platform/metal/metal_data_types.h"
+#include "taichi/backends/metal/data_types.h"
 #include "taichi/program/kernel.h"
 
 // Data structures defined in this file may overlap with some of the Taichi data

diff --git a/taichi/platform/metal/metal_runtime.cpp → taichi/backends/metal/runtime.cpp b/taichi/platform/metal/metal_runtime.cpp → taichi/backends/metal/runtime.cpp
@@ -1,4 +1,4 @@
-#include "taichi/platform/metal/metal_runtime.h"
+#include "taichi/backends/metal/runtime.h"
 
 #include <algorithm>
 #include <cstring>
@@ -12,7 +12,7 @@
 #include <sys/mman.h>
 #include <unistd.h>
 
-#include "taichi/platform/metal/metal_api.h"
+#include "taichi/backends/metal/api.h"
 #endif  // TI_PLATFORM_OSX
 
 TLANG_NAMESPACE_BEGIN
@@ -22,7 +22,7 @@ namespace metal {
 
 namespace {
 namespace shaders {
-#include "taichi/platform/metal/shaders/runtime_utils.metal.h"
+#include "taichi/backends/metal/shaders/runtime_utils.metal.h"
 }
 
 using KernelTaskType = OffloadedStmt::TaskType;
@@ -189,7 +189,7 @@ class RuntimeListOpsMtlKernel : public CompiledMtlKernelBase {
   // args[1] = child_snode_id
   // Note that this args buffer has nothing to do with the one passed to Taichi
   // kernel.
-  // See taichi/platform/metal/shaders/runtime_kernels.metal.h
+  // See taichi/backends/metal/shaders/runtime_kernels.metal.h
   std::unique_ptr<BufferMemoryView> args_mem_;
   nsobj_unique_ptr<MTLBuffer> args_buffer_;
 };

diff --git a/taichi/platform/metal/metal_runtime.h → taichi/backends/metal/runtime.h b/taichi/platform/metal/metal_runtime.h → taichi/backends/metal/runtime.h
@@ -5,7 +5,7 @@
 #include <unordered_map>
 
 #include "taichi/lang_util.h"
-#include "taichi/platform/metal/metal_kernel_util.h"
+#include "taichi/backends/metal/kernel_util.h"
 #include "taichi/program/profiler.h"
 #include "taichi/struct/struct_metal.h"
 #include "taichi/system/memory_pool.h"

diff --git a/taichi/platform/metal/shaders/atomic_stubs.h → taichi/backends/metal/shaders/atomic_stubs.h b/taichi/platform/metal/shaders/atomic_stubs.h → taichi/backends/metal/shaders/atomic_stubs.h
diff --git a/taichi/platform/metal/shaders/epilog.h → taichi/backends/metal/shaders/epilog.h b/taichi/platform/metal/shaders/epilog.h → taichi/backends/metal/shaders/epilog.h
diff --git a/...hi/platform/metal/shaders/helpers.metal.h → ...hi/backends/metal/shaders/helpers.metal.h b/...hi/platform/metal/shaders/helpers.metal.h → ...hi/backends/metal/shaders/helpers.metal.h
@@ -1,4 +1,4 @@
-#include "taichi/platform/metal/shaders/prolog.h"
+#include "taichi/backends/metal/shaders/prolog.h"
 
 #ifdef TI_INSIDE_METAL_CODEGEN
 
@@ -19,14 +19,18 @@ static_assert(false, "Do not include");
 
 #endif  // TI_INSIDE_METAL_CODEGEN
 
+// clang-format off
 METAL_BEGIN_HELPERS_DEF
-STR(template <typename T, typename G> T union_cast(G g) {
-  // For some reason, if I emit taichi/common.h's union_cast(), Metal failed
-  // to compile. More strangely, if I copy the generated code to XCode as a
-  // Metal kernel, it compiled successfully...
-  static_assert(sizeof(T) == sizeof(G), "Size mismatch");
-  return *reinterpret_cast<thread const T *>(&g);
-}
+STR(
+    // clang-format on
+    template <typename T, typename G>
+    T union_cast(G g) {
+      // For some reason, if I emit taichi/common.h's union_cast(), Metal failed
+      // to compile. More strangely, if I copy the generated code to XCode as a
+      // Metal kernel, it compiled successfully...
+      static_assert(sizeof(T) == sizeof(G), "Size mismatch");
+      return *reinterpret_cast<thread const T *>(&g);
+    }
 
     inline int ifloordiv(int lhs, int rhs) {
       const int intm = (lhs / rhs);
@@ -87,10 +91,13 @@ STR(template <typename T, typename G> T union_cast(G g) {
             metal::memory_order_relaxed);
       }
       return old_val;
-    })
+    }
+    // clang-format off
+)
 METAL_END_HELPERS_DEF
+// clang-format on
 
 #undef METAL_BEGIN_HELPERS_DEF
 #undef METAL_END_HELPERS_DEF
 
-#include "taichi/platform/metal/shaders/epilog.h"
+#include "taichi/backends/metal/shaders/epilog.h"
diff --git a/taichi/platform/metal/shaders/prolog.h → taichi/backends/metal/shaders/prolog.h b/taichi/platform/metal/shaders/prolog.h → taichi/backends/metal/shaders/prolog.h
@@ -24,6 +24,6 @@
 
 #define byte char
 
-#include "taichi/platform/metal/shaders/atomic_stubs.h"
+#include "taichi/backends/metal/shaders/atomic_stubs.h"
 
 #endif  // TI_INSIDE_METAL_CODEGEN
diff --git a/...orm/metal/shaders/runtime_kernels.metal.h → ...nds/metal/shaders/runtime_kernels.metal.h b/...orm/metal/shaders/runtime_kernels.metal.h → ...nds/metal/shaders/runtime_kernels.metal.h
@@ -1,14 +1,14 @@
 #ifndef TI_METAL_NESTED_INCLUDE
 
 #define TI_METAL_NESTED_INCLUDE
-#include "taichi/platform/metal/shaders/runtime_utils.metal.h"
+#include "taichi/backends/metal/shaders/runtime_utils.metal.h"
 #undef TI_METAL_NESTED_INCLUDE
 
 #else
-#include "taichi/platform/metal/shaders/runtime_utils.metal.h"
+#include "taichi/backends/metal/shaders/runtime_utils.metal.h"
 #endif  // TI_METAL_NESTED_INCLUDE
 
-#include "taichi/platform/metal/shaders/prolog.h"
+#include "taichi/backends/metal/shaders/prolog.h"
 
 #ifdef TI_INSIDE_METAL_CODEGEN
 
@@ -39,23 +39,26 @@ struct Runtime {
 
 #endif  // TI_INSIDE_METAL_CODEGEN
 
+// clang-format off
 METAL_BEGIN_RUNTIME_KERNELS_DEF
-STR(kernel void clear_list(device byte *runtime_addr[[buffer(0)]],
-                           device int *args[[buffer(1)]],
-                           const uint utid_[[thread_position_in_grid]]) {
-  if (utid_ > 0)
-    return;
-  int child_snode_id = args[1];
-  device ListManager *child_list =
-      &(reinterpret_cast<device Runtime *>(runtime_addr)
-            ->snode_lists[child_snode_id]);
-  clear(child_list);
-}
+STR(
+    // clang-format on
+    kernel void clear_list(device byte *runtime_addr [[buffer(0)]],
+                           device int *args [[buffer(1)]],
+                           const uint utid_ [[thread_position_in_grid]]) {
+      if (utid_ > 0)
+        return;
+      int child_snode_id = args[1];
+      device ListManager *child_list =
+          &(reinterpret_cast<device Runtime *>(runtime_addr)
+                ->snode_lists[child_snode_id]);
+      clear(child_list);
+    }
 
-    kernel void element_listgen(device byte *runtime_addr[[buffer(0)]],
-                                device byte *root_addr[[buffer(1)]],
-                                device int *args[[buffer(2)]],
-                                const uint utid_[[thread_position_in_grid]]) {
+    kernel void element_listgen(device byte *runtime_addr [[buffer(0)]],
+                                device byte *root_addr [[buffer(1)]],
+                                device int *args [[buffer(2)]],
+                                const uint utid_ [[thread_position_in_grid]]) {
       device Runtime *runtime =
           reinterpret_cast<device Runtime *>(runtime_addr);
       device byte *list_data_addr =
@@ -86,10 +89,13 @@ STR(kernel void clear_list(device byte *runtime_addr[[buffer(0)]],
           append(child_list, child_elem, list_data_addr);
         }
       }
-    })
+    }
+    // clang-format off
+)
 METAL_END_RUNTIME_KERNELS_DEF
+// clang-format on
 
 #undef METAL_BEGIN_RUNTIME_KERNELS_DEF
 #undef METAL_END_RUNTIME_KERNELS_DEF
 
-#include "taichi/platform/metal/shaders/epilog.h"
+#include "taichi/backends/metal/shaders/epilog.h"
diff --git a/...orm/metal/shaders/runtime_structs.metal.h → ...nds/metal/shaders/runtime_structs.metal.h b/...orm/metal/shaders/runtime_structs.metal.h → ...nds/metal/shaders/runtime_structs.metal.h
@@ -1,4 +1,4 @@
-#include "taichi/platform/metal/shaders/prolog.h"
+#include "taichi/backends/metal/shaders/prolog.h"
 
 #ifdef TI_INSIDE_METAL_CODEGEN
 
@@ -25,8 +25,11 @@ static_assert(taichi_max_num_indices == 8,
 
 #endif  // TI_INSIDE_METAL_CODEGEN
 
+// clang-format off
 METAL_BEGIN_RUNTIME_STRUCTS_DEF
-STR(constant constexpr int kTaichiMaxNumIndices = 8;
+STR(
+    // clang-format on
+    constant constexpr int kTaichiMaxNumIndices = 8;
 
     struct ListgenElement {
       int32_t coords[kTaichiMaxNumIndices];
@@ -61,10 +64,13 @@ STR(constant constexpr int kTaichiMaxNumIndices = 8;
       };
 
       Extractor extractors[kTaichiMaxNumIndices];
-    };)
+    };
+    // clang-format off
+)
 METAL_END_RUNTIME_STRUCTS_DEF
+// clang-format on
 
 #undef METAL_BEGIN_RUNTIME_STRUCTS_DEF
 #undef METAL_END_RUNTIME_STRUCTS_DEF
 
-#include "taichi/platform/metal/shaders/epilog.h"
+#include "taichi/backends/metal/shaders/epilog.h"
diff --git a/...tform/metal/shaders/runtime_utils.metal.h → ...kends/metal/shaders/runtime_utils.metal.h b/...tform/metal/shaders/runtime_utils.metal.h → ...kends/metal/shaders/runtime_utils.metal.h
@@ -1,14 +1,14 @@
 #ifndef TI_METAL_NESTED_INCLUDE
 
 #define TI_METAL_NESTED_INCLUDE
-#include "taichi/platform/metal/shaders/runtime_structs.metal.h"
+#include "taichi/backends/metal/shaders/runtime_structs.metal.h"
 #undef TI_METAL_NESTED_INCLUDE
 
 #else
-#include "taichi/platform/metal/shaders/runtime_structs.metal.h"
+#include "taichi/backends/metal/shaders/runtime_structs.metal.h"
 #endif  // TI_METAL_NESTED_INCLUDE
 
-#include "taichi/platform/metal/shaders/prolog.h"
+#include "taichi/backends/metal/shaders/prolog.h"
 
 #ifdef TI_INSIDE_METAL_CODEGEN
 
@@ -112,11 +112,12 @@ STR(
         const int addition = (((l >> ex.acc_offset) & mask) << ex.start);
         child_elem->coords[i] = (parent_elem.coords[i] | addition);
       }
-    })
+    }
+)
 METAL_END_RUNTIME_UTILS_DEF
 // clang-format on
 
 #undef METAL_BEGIN_RUNTIME_UTILS_DEF
 #undef METAL_END_RUNTIME_UTILS_DEF
 
-#include "taichi/platform/metal/shaders/epilog.h"
+#include "taichi/backends/metal/shaders/epilog.h"
diff --git a/taichi/codegen/codegen_metal.cpp b/taichi/codegen/codegen_metal.cpp
@@ -11,7 +11,7 @@ namespace metal {
 namespace {
 
 namespace shaders {
-#include "taichi/platform/metal/shaders/runtime_structs.metal.h"
+#include "taichi/backends/metal/shaders/runtime_structs.metal.h"
 }  // namespace shaders
 
 using BuffersEnum = MetalKernelAttributes::Buffers;
@@ -488,7 +488,7 @@ class MetalKernelCodegen : public IRVisitor {
     emit("using byte = uchar;");
     emit("");
 #define TI_INSIDE_METAL_CODEGEN
-#include "taichi/platform/metal/shaders/helpers.metal.h"
+#include "taichi/backends/metal/shaders/helpers.metal.h"
     line_appender_.append_raw(kMetalHelpersSourceCode);
 #undef TI_INSIDE_METAL_CODEGEN
     emit("");

diff --git a/taichi/codegen/codegen_metal.h b/taichi/codegen/codegen_metal.h
@@ -7,9 +7,9 @@
 
 #include "taichi/inc/constants.h"
 #include "taichi/lang_util.h"
-#include "taichi/platform/metal/metal_data_types.h"
-#include "taichi/platform/metal/metal_kernel_util.h"
-#include "taichi/platform/metal/metal_runtime.h"
+#include "taichi/backends/metal/data_types.h"
+#include "taichi/backends/metal/kernel_util.h"
+#include "taichi/backends/metal/runtime.h"
 #include "taichi/program/program.h"
 #include "taichi/struct/struct_metal.h"
 

diff --git a/taichi/program/program.cpp b/taichi/program/program.cpp
@@ -2,10 +2,9 @@
 
 #include "program.h"
 
-#include <taichi/common/task.h>
-#include <taichi/platform/metal/metal_api.h>
-#include <taichi/platform/opengl/opengl_api.h>
-
+#include "taichi/common/task.h"
+#include "taichi/backends/metal/api.h"
+#include "taichi/platform/opengl/opengl_api.h"
 #include "taichi/codegen/codegen_cuda.h"
 #include "taichi/codegen/codegen_metal.h"
 #include "taichi/codegen/codegen_opengl.h"

diff --git a/taichi/program/program.h b/taichi/program/program.h
@@ -10,7 +10,7 @@
 #include "taichi/ir/snode.h"
 #include "taichi/lang_util.h"
 #include "taichi/llvm/llvm_context.h"
-#include "taichi/platform/metal/metal_runtime.h"
+#include "taichi/backends/metal/runtime.h"
 #include "taichi/platform/opengl/opengl_kernel_util.h"
 #include "taichi/program/kernel.h"
 #include "taichi/program/profiler.h"

diff --git a/taichi/python/export_misc.cpp b/taichi/python/export_misc.cpp
@@ -12,7 +12,7 @@
 #include <taichi/system/profiler.h>
 #include <taichi/system/memory_usage_monitor.h>
 #include <taichi/system/unit_dll.h>
-#include <taichi/platform/metal/metal_api.h>
+#include <taichi/backends/metal/api.h>
 #include <taichi/platform/opengl/opengl_api.h>
 #if defined(TI_WITH_CUDA)
 #include <cuda_runtime_api.h>

diff --git a/taichi/struct/struct_metal.cpp b/taichi/struct/struct_metal.cpp
@@ -7,21 +7,21 @@
 #include <vector>
 
 #include "taichi/math/arithmetic.h"
-#include "taichi/platform/metal/metal_data_types.h"
-#include "taichi/platform/metal/metal_kernel_util.h"
+#include "taichi/backends/metal/data_types.h"
+#include "taichi/backends/metal/kernel_util.h"
 #include "taichi/util/line_appender.h"
 
 TLANG_NAMESPACE_BEGIN
 namespace metal {
 namespace {
 namespace shaders {
 #define TI_INSIDE_METAL_CODEGEN
-#include "taichi/platform/metal/shaders/runtime_kernels.metal.h"
-#include "taichi/platform/metal/shaders/runtime_structs.metal.h"
-#include "taichi/platform/metal/shaders/runtime_utils.metal.h"
+#include "taichi/backends/metal/shaders/runtime_kernels.metal.h"
+#include "taichi/backends/metal/shaders/runtime_structs.metal.h"
+#include "taichi/backends/metal/shaders/runtime_utils.metal.h"
 #undef TI_INSIDE_METAL_CODEGEN
 
-#include "taichi/platform/metal/shaders/runtime_structs.metal.h"
+#include "taichi/backends/metal/shaders/runtime_structs.metal.h"
 
 }  // namespace shaders