Resolve merge conflicts

.clang-format

@@ -92,4 +92,6 @@ SpacesInSquareBrackets: false
 Standard:        Auto
 TabWidth:        8
 UseTab:          Never
+---
+Language: ObjC
 ...

.github/workflows/release.yml

@@ -274,6 +274,11 @@ jobs:
       matrix: ${{ fromJson(needs.matrix_prep.outputs.matrix) }}
     runs-on: [self-hosted, windows, cuda, OpenGL]
     steps:
+      - name: Workaround checkout Needed single revision issue
+        run: |
+          git config --system core.longpaths true
+          git submodule foreach 'git rev-parse HEAD > /dev/null 2>&1 || rm -rf $PWD'
+
       - uses: actions/checkout@v3
         with:
           submodules: 'recursive'

.github/workflows/scripts/aot-demo.sh

@@ -4,7 +4,8 @@ set -ex
 export TI_SKIP_VERSION_CHECK=ON
 export TI_CI=1
 
-export TAICHI_AOT_DEMO_URL=https://github.com/bobcao3/taichi-aot-demo
+# IF YOU PIN THIS TO A COMMIT/BRANCH, YOU'RE RESPONSIBLE TO REVERT IT BACK TO MASTER ONCE MERGED.
+export TAICHI_AOT_DEMO_URL=https://github.com/taichi-dev/taichi-aot-demo
 export TAICHI_AOT_DEMO_BRANCH=master
 
 export TAICHI_UNITY2_URL=https://github.com/taichi-dev/taichi-unity2
@@ -27,7 +28,6 @@ function build-and-smoke-test-android-aot-demo {
     export TAICHI_REPO_DIR=$(pwd)/taichi
 
     rm -rf taichi-aot-demo
-    # IF YOU PIN THIS TO A COMMIT/BRANCH, YOU'RE RESPONSIBLE TO REVERT IT BACK TO MASTER ONCE MERGED.
     git clone --recursive --depth=1 -b "$TAICHI_AOT_DEMO_BRANCH" "$TAICHI_AOT_DEMO_URL"
 
     # Install taichi-python

.github/workflows/scripts/build.py

@@ -28,10 +28,6 @@ def setup_llvm(env_out: dict) -> None:
         if 'AMDGPU_TEST' in os.environ:
             # FIXME: AMDGPU bots are currently maintained separately,
             #        we should unify them with the rest of the bots.
-            lnsf = sh.sudo.ln.bake('-sf')
-            lnsf('/usr/bin/clang++-10', '/usr/bin/clang++')
-            lnsf('/usr/bin/clang-10', '/usr/bin/clang')
-            lnsf('/usr/bin/ld.lld-10', '/usr/bin/ld.lld')
             env_out['LLVM_DIR'] = '/taichi-llvm-15'
             return
         elif is_manylinux2014():

.github/workflows/testing.yml

@@ -291,7 +291,7 @@ jobs:
           . .github/workflows/scripts/common-utils.sh
 
           ci-docker-run-amdgpu --name taichi-build \
-            registry.taichigraphics.com/taichidev-ubuntu18.04.amdgpu:v0.0.5 \
+            registry.taichigraphics.com/taichidev-ubuntu18.04.amdgpu:v0.0.6 \
             /home/dev/taichi/.github/workflows/scripts/build.py
 
         env:
@@ -311,7 +311,7 @@ jobs:
           . .github/workflows/scripts/common-utils.sh
 
           ci-docker-run-amdgpu --name taichi-test \
-             registry.taichigraphics.com/taichidev-ubuntu18.04.amdgpu:v0.0.5 \
+             registry.taichigraphics.com/taichidev-ubuntu18.04.amdgpu:v0.0.6 \
              /home/dev/taichi/.github/workflows/scripts/unix_test.sh
         env:
           PY: '3.8'
@@ -351,6 +351,7 @@ jobs:
     steps:
       - name: Workaround checkout Needed single revision issue
         run: |
+          git config --system core.longpaths true
           git submodule foreach 'git rev-parse HEAD > /dev/null 2>&1 || rm -rf $PWD'
 
       - uses: actions/checkout@v3

CMakeLists.txt

@@ -16,6 +16,7 @@ if (NOT DEFINED TI_VERSION_MAJOR)
 endif()
 
 set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_EXPORT_COMPILECOMMANDS ON)
 
 execute_process(
   WORKING_DIRECTORY ${CMAKE_SOURCE_DIR}

c_api/include/taichi/cpp/taichi.hpp

@@ -771,8 +771,8 @@ class CapabilityLevelConfigBuilder {
         value ? TI_TRUE : TI_FALSE;
     return *this;
   }
-  Self &spirv_has_atomic_i64(bool value = true) {
-    cap_level_infos[TI_CAPABILITY_SPIRV_HAS_ATOMIC_I64] =
+  Self &spirv_has_atomic_int64(bool value = true) {
+    cap_level_infos[TI_CAPABILITY_SPIRV_HAS_ATOMIC_INT64] =
         value ? TI_TRUE : TI_FALSE;
     return *this;
   }

c_api/include/taichi/taichi_core.h

@@ -386,7 +386,7 @@ typedef enum TiCapability {
   TI_CAPABILITY_SPIRV_HAS_INT64 = 4,
   TI_CAPABILITY_SPIRV_HAS_FLOAT16 = 5,
   TI_CAPABILITY_SPIRV_HAS_FLOAT64 = 6,
-  TI_CAPABILITY_SPIRV_HAS_ATOMIC_I64 = 7,
+  TI_CAPABILITY_SPIRV_HAS_ATOMIC_INT64 = 7,
   TI_CAPABILITY_SPIRV_HAS_ATOMIC_FLOAT16 = 8,
   TI_CAPABILITY_SPIRV_HAS_ATOMIC_FLOAT16_ADD = 9,
   TI_CAPABILITY_SPIRV_HAS_ATOMIC_FLOAT16_MINMAX = 10,

c_api/src/taichi_core_impl.cpp

@@ -2,6 +2,7 @@
 #include "taichi_opengl_impl.h"
 #include "taichi_vulkan_impl.h"
 #include "taichi_llvm_impl.h"
+#include "taichi_metal_impl.h"
 #include "taichi/program/ndarray.h"
 #include "taichi/program/texture.h"
 #include "taichi/common/virtual_dir.h"
@@ -49,6 +50,14 @@ bool is_arm64_available() {
 #endif
 }
 
+bool is_metal_available() {
+#if defined(__APPLE__) && defined(TI_WITH_METAL)
+  return taichi::lang::metal::is_metal_api_available();
+#else
+  return false;
+#endif  // __APPLE__
+}
+
 struct ErrorCache {
   TiError error{TI_ERROR_SUCCESS};
   std::string message{};
@@ -167,6 +176,9 @@ void ti_get_available_archs(uint32_t *arch_count, TiArch *archs) {
     if (is_arm64_available()) {
       AVAILABLE_ARCHS.emplace_back(TI_ARCH_ARM64);
     }
+    if (is_metal_available()) {
+      AVAILABLE_ARCHS.emplace_back(TI_ARCH_METAL);
+    }
   }
 
   size_t n = std::min((size_t)*arch_count, AVAILABLE_ARCHS.size());
@@ -257,6 +269,13 @@ TiRuntime ti_create_runtime(TiArch arch, uint32_t device_index) {
       break;
     }
 #endif  // TI_WITH_LLVM
+#ifdef TI_WITH_METAL
+    case TI_ARCH_METAL: {
+      out = (TiRuntime)(static_cast<Runtime *>(
+          new capi::MetalRuntime(taichi::Arch::metal)));
+      break;
+    }
+#endif  // TI_WITH_METAL
     default: {
       TI_CAPI_NOT_SUPPORTED(arch);
       return TI_NULL_HANDLE;

c_api/src/taichi_core_impl.h

@@ -173,6 +173,7 @@ class Runtime {
   virtual void wait() = 0;
 
   class VulkanRuntime *as_vk();
+  class MetalRuntime *as_mtl();
 };
 
 class AotModule {

c_api/src/taichi_metal_impl.h

@@ -0,0 +1,27 @@
+#pragma once
+#ifdef TI_WITH_METAL
+#include "taichi_core_impl.h"
+#include "taichi_gfx_impl.h"
+#include "taichi/rhi/metal/metal_device.h"
+
+namespace capi {
+
+class MetalRuntime;
+
+class MetalRuntime : public GfxRuntime {
+ private:
+  std::unique_ptr<taichi::lang::metal::MetalDevice> mtl_device_;
+  taichi::lang::gfx::GfxRuntime gfx_runtime_;
+
+ public:
+  explicit MetalRuntime(taichi::Arch arch);
+
+  taichi::lang::Device &get() override;
+  taichi::lang::gfx::GfxRuntime &get_gfx_runtime() override;
+
+  taichi::lang::metal::MetalDevice &get_mtl();
+};
+
+}  // namespace capi
+
+#endif  // TI_WITH_METAL

c_api/src/taichi_metal_impl.mm

@@ -0,0 +1,26 @@
+#ifdef TI_WITH_METAL
+#include "taichi_metal_impl.h"
+#include "taichi/rhi/metal/metal_device.h"
+#include "taichi/runtime/gfx/runtime.h"
+
+namespace capi {
+
+MetalRuntime::MetalRuntime(taichi::Arch arch)
+    : GfxRuntime(arch), mtl_device_(taichi::lang::metal::MetalDevice::create()),
+      gfx_runtime_(taichi::lang::gfx::GfxRuntime::Params{
+          host_result_buffer_.data(), mtl_device_.get()}) {}
+
+taichi::lang::Device &MetalRuntime::get() {
+  return static_cast<taichi::lang::Device &>(*mtl_device_);
+}
+taichi::lang::gfx::GfxRuntime &MetalRuntime::get_gfx_runtime() {
+  return gfx_runtime_;
+}
+
+taichi::lang::metal::MetalDevice &MetalRuntime::get_mtl() {
+  return *mtl_device_;
+}
+
+} // namespace capi
+
+#endif // TI_WITH_METAL

c_api/tests/c_api_behavior_test.cpp

@@ -20,7 +20,6 @@ TEST_F(CapiTest, TestBehaviorCreateRuntime) {
   inner(TI_ARCH_JS);
   inner(TI_ARCH_CC);
   inner(TI_ARCH_WASM);
-  inner(TI_ARCH_METAL);
   inner(TI_ARCH_DX11);
   inner(TI_ARCH_DX12);
   inner(TI_ARCH_OPENCL);

c_api/tests/c_api_cgraph_test.cpp

@@ -19,23 +19,34 @@ void graph_aot_test(TiArch arch) {
   ti::AotModule aot_mod = runtime.load_aot_module(aot_mod_ss.str().c_str());
   ti::ComputeGraph run_graph = aot_mod.get_compute_graph("run_graph");
 
-  ti::NdArray<int32_t> arr_array =
+  ti::NdArray<int32_t> arr_array_0 =
+      runtime.allocate_ndarray<int32_t>({kArrLen}, {}, true);
+  ti::NdArray<int32_t> arr_array_1 =
       runtime.allocate_ndarray<int32_t>({kArrLen}, {1}, true);
 
   run_graph["base0"] = base0_val;
   run_graph["base1"] = base1_val;
   run_graph["base2"] = base2_val;
-  run_graph["arr"] = arr_array;
+  run_graph["arr0"] = arr_array_0;
+  run_graph["arr1"] = arr_array_1;
   run_graph.launch();
   runtime.wait();
 
   // Check Results
-  auto *data = reinterpret_cast<int32_t *>(arr_array.map());
+  auto *data = reinterpret_cast<int32_t *>(arr_array_0.map());
 
   for (int i = 0; i < kArrLen; i++) {
     EXPECT_EQ(data[i], 3 * i + base0_val + base1_val + base2_val);
   }
-  arr_array.unmap();
+
+  data = reinterpret_cast<int32_t *>(arr_array_1.map());
+
+  for (int i = 0; i < kArrLen; i++) {
+    EXPECT_EQ(data[i], 3 * i + base0_val + base1_val + base2_val);
+  }
+
+  arr_array_0.unmap();
+  arr_array_1.unmap();
 }
 
 void texture_aot_test(TiArch arch) {

ci/Dockerfile.ubuntu.18.04.amdgpu

@@ -11,6 +11,7 @@ RUN apt-get update && \
                        libtinfo-dev \
                        sudo \
                        clang-10 \
+                       lld-10 \
                        wget \
                        git \
                        unzip \
@@ -44,10 +45,14 @@ RUN apt-get update && \
     printf "root ALL=(ALL:ALL) NOPASSWD: ALL\ndev ALL=(ALL:ALL) NOPASSWD: ALL\n" > /etc/sudoers && \
     true
 
+RUN ln -sf /usr/bin/clang++-10 /usr/bin/clang++ && \
+    ln -sf /usr/bin/clang-10 /usr/bin/clang && \
+    ln -sf /usr/bin/ld.lld-10 /usr/bin/ld.lld
+
 # Install LLVM 15
 WORKDIR /
 # Make sure this URL gets updated each time there is a new prebuilt bin release
-RUN wget https://github.com/GaleSeLee/assets/releases/download/v0.0.2/taichi-llvm-15.0.0-linux.zip
+RUN wget https://github.com/GaleSeLee/assets/releases/download/v0.0.4/taichi-llvm-15.0.0-linux.zip
 RUN unzip taichi-llvm-15.0.0-linux.zip && \
     rm taichi-llvm-15.0.0-linux.zip
 ENV PATH="/taichi-llvm-15.0.0-linux/bin:$PATH"

cmake/TaichiCAPI.cmake

@@ -26,14 +26,18 @@ if (TI_WITH_LLVM)
   list(APPEND C_API_SOURCE "c_api/src/taichi_llvm_impl.cpp")
 endif()
 
-if (TI_WITH_OPENGL OR TI_WITH_VULKAN)
+if (TI_WITH_OPENGL OR TI_WITH_VULKAN OR TI_WITH_METAL)
   list(APPEND C_API_SOURCE "c_api/src/taichi_gfx_impl.cpp")
 endif()
 
 if (TI_WITH_OPENGL)
   list(APPEND C_API_SOURCE "c_api/src/taichi_opengl_impl.cpp")
 endif()
 
+if (TI_WITH_METAL)
+  list(APPEND C_API_SOURCE "c_api/src/taichi_metal_impl.mm")
+endif()
+
 if (TI_WITH_VULKAN)
   list(APPEND C_API_SOURCE "c_api/src/taichi_vulkan_impl.cpp")
   if (APPLE)

cmake/TaichiCore.cmake

@@ -61,6 +61,11 @@ if (APPLE)
         set(TI_WITH_AMDGPU OFF)
         message(WARNING "AMDGPU backend not supported on OS X. Setting TI_WITH_AMDGPU to OFF.")
     endif()
+else()
+    if (TI_WITH_METAL)
+        set(TI_WITH_METAL OFF)
+        message(WARNING "Metal backend only supported on OS X. Setting TI_WITH_METAL to OFF.")
+    endif()
 endif()
 
 if (WIN32)
@@ -236,6 +241,9 @@ if(TI_WITH_LLVM)
     if (TI_WITH_AMDGPU)
         llvm_map_components_to_libnames(llvm_amdgpu_libs AMDGPU)
         add_subdirectory(taichi/rhi/amdgpu)
+        add_subdirectory(taichi/codegen/amdgpu)
+
+        target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE amdgpu_codegen)
         target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE amdgpu_rhi)
     endif()
 
@@ -296,13 +304,7 @@ if (TI_WITH_METAL)
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_WITH_METAL")
 
     add_subdirectory(taichi/rhi/metal)
-    add_subdirectory(taichi/runtime/metal)
     add_subdirectory(taichi/runtime/program_impls/metal)
-    add_subdirectory(taichi/codegen/metal)
-    add_subdirectory(taichi/cache/metal)
-
-    target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE metal_codegen)
-    target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE metal_runtime)
     target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE metal_program_impl)
 endif()
 
@@ -331,12 +333,12 @@ add_subdirectory(taichi/codegen/spirv)
 add_subdirectory(taichi/cache/gfx)
 add_subdirectory(taichi/runtime/gfx)
 
-if (TI_WITH_OPENGL OR TI_WITH_VULKAN OR TI_WITH_DX11)
+if (TI_WITH_OPENGL OR TI_WITH_VULKAN OR TI_WITH_DX11 OR TI_WITH_METAL)
   target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE spirv_codegen)
   target_link_libraries(${CORE_LIBRARY_NAME} PRIVATE gfx_runtime)
 endif()
 
-if (TI_WITH_OPENGL OR TI_WITH_DX11)
+if (TI_WITH_OPENGL OR TI_WITH_DX11 OR TI_WITH_METAL)
   set(SPIRV_CROSS_CLI false)
   add_subdirectory(${PROJECT_SOURCE_DIR}/external/SPIRV-Cross ${PROJECT_BINARY_DIR}/external/SPIRV-Cross)
 endif()

cmake/TaichiExamples.cmake

@@ -25,12 +25,10 @@ target_link_libraries(${EXAMPLES_NAME} PRIVATE taichi_core)
 if (TI_WITH_METAL)
   target_link_libraries(${EXAMPLES_NAME} PRIVATE
     metal_program_impl
-    metal_runtime
-    metal_codegen
   )
 endif()
 
-if (TI_WITH_VULKAN OR TI_WITH_OPENGL)
+if (TI_WITH_VULKAN OR TI_WITH_OPENGL OR TI_WITH_METAL)
   target_link_libraries(${EXAMPLES_NAME} PRIVATE gfx_runtime)
 endif()
 
@@ -42,6 +40,10 @@ if (TI_WITH_OPENGL)
   target_link_libraries(${EXAMPLES_NAME} PRIVATE opengl_rhi)
 endif()
 
+if (TI_WITH_METAL)
+  target_link_libraries(${EXAMPLES_NAME} PRIVATE metal_rhi)
+endif()
+
 # TODO 4832: be specific on the header dependencies here, e.g., ir
 target_include_directories(${EXAMPLES_NAME}
   PRIVATE

docs/lang/articles/c-api/taichi_core.md

@@ -409,7 +409,7 @@ typedef enum TiCapability {
   TI_CAPABILITY_SPIRV_HAS_INT64 = 4,
   TI_CAPABILITY_SPIRV_HAS_FLOAT16 = 5,
   TI_CAPABILITY_SPIRV_HAS_FLOAT64 = 6,
-  TI_CAPABILITY_SPIRV_HAS_ATOMIC_I64 = 7,
+  TI_CAPABILITY_SPIRV_HAS_ATOMIC_INT64 = 7,
   TI_CAPABILITY_SPIRV_HAS_ATOMIC_FLOAT16 = 8,
   TI_CAPABILITY_SPIRV_HAS_ATOMIC_FLOAT16_ADD = 9,
   TI_CAPABILITY_SPIRV_HAS_ATOMIC_FLOAT16_MINMAX = 10,

docs/lang/articles/glossary/_category_.json

@@ -0,0 +1,4 @@
+{
+  "label": "Glossary",
+  "position": 17
+}