Add native Apple Metal/MPS backend

.github/workflows/ci.yml

@@ -231,6 +231,147 @@ jobs:
         run: |
           tests/ctranslate2_test tests/data
 
+  build-and-test-mps:
+    runs-on: macos-15
+    timeout-minutes: 30
+    env:
+      CT2_VERBOSE: 1
+    defaults:
+      run:
+        shell: bash
+
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          submodules: recursive
+
+      - name: Verify Apple Silicon and execute a Metal kernel
+        run: |
+          test "$(uname -m)" = "arm64"
+          system_profiler SPHardwareDataType SPDisplaysDataType | tee runner-hardware.txt
+          cat > "$RUNNER_TEMP/metal-probe.swift" <<'SWIFT'
+          import Foundation
+          import Metal
+
+          guard let device = MTLCreateSystemDefaultDevice() else {
+              fatalError("No Metal device available on this runner")
+          }
+          print("Metal device: \(device.name)")
+          print("Unified memory: \(device.hasUnifiedMemory)")
+          print("Recommended working set: \(device.recommendedMaxWorkingSetSize) bytes")
+          let source = """
+          #include <metal_stdlib>
+          using namespace metal;
+          kernel void increment(device uint *values [[buffer(0)]],
+                                uint index [[thread_position_in_grid]]) {
+              values[index] += 1;
+          }
+          """
+          let library = try device.makeLibrary(source: source, options: nil)
+          guard let function = library.makeFunction(name: "increment"),
+                let queue = device.makeCommandQueue(),
+                let buffer = device.makeBuffer(length: 4 * MemoryLayout<UInt32>.stride,
+                                               options: .storageModeShared),
+                let commandBuffer = queue.makeCommandBuffer(),
+                let encoder = commandBuffer.makeComputeCommandEncoder() else {
+              fatalError("Could not create Metal compute resources")
+          }
+          let pipeline = try device.makeComputePipelineState(function: function)
+          let values = buffer.contents().bindMemory(to: UInt32.self, capacity: 4)
+          for index in 0..<4 { values[index] = UInt32(index) }
+          encoder.setComputePipelineState(pipeline)
+          encoder.setBuffer(buffer, offset: 0, index: 0)
+          encoder.dispatchThreads(MTLSize(width: 4, height: 1, depth: 1),
+                                  threadsPerThreadgroup: MTLSize(width: 4, height: 1, depth: 1))
+          encoder.endEncoding()
+          commandBuffer.commit()
+          commandBuffer.waitUntilCompleted()
+          guard commandBuffer.status == .completed else {
+              fatalError("Metal command failed: \(String(describing: commandBuffer.error))")
+          }
+          for index in 0..<4 {
+              precondition(values[index] == UInt32(index + 1), "Incorrect Metal kernel result")
+          }
+          print("Metal compute probe passed")
+          SWIFT
+          swift "$RUNNER_TEMP/metal-probe.swift" | tee metal-probe.txt
+
+      - name: Set up Python 3.12
+        uses: actions/setup-python@v5
+        with:
+          python-version: "3.12"
+
+      - name: Configure MPS build
+        run: |
+          cmake -S . -B build-mps \
+            -DCMAKE_BUILD_TYPE=Release \
+            -DCMAKE_INSTALL_PREFIX="$PWD/install-mps" \
+            -DCMAKE_OSX_ARCHITECTURES=arm64 \
+            -DCMAKE_POLICY_VERSION_MINIMUM=3.5 \
+            -DBUILD_TESTS=ON \
+            -DWITH_MPS=ON \
+            -DWITH_ACCELERATE=ON \
+            -DWITH_MKL=OFF \
+            -DWITH_RUY=ON \
+            -DOPENMP_RUNTIME=NONE
+
+      - name: Build and install MPS library
+        run: |
+          cmake --build build-mps --parallel "$(sysctl -n hw.logicalcpu)"
+          cmake --install build-mps
+
+      - name: Download test data
+        id: test_data
+        working-directory: tests/data/models
+        run: |
+          curl --fail --location --output transliteration-aren-all.tar.gz \
+            https://opennmt-models.s3.amazonaws.com/transliteration-aren-all.tar.gz
+          tar xf transliteration-aren-all.tar.gz
+
+      - name: Run MPS runtime test suite
+        run: |
+          build-mps/tests/ctranslate2_test tests/data \
+            --gtest_filter='MPSBackendTest.*:MPSPrimitiveTest.*:TranslatorTest.MPS*:MPS/*' \
+            --gtest_output=xml:mps-results.xml
+
+      - name: Run CPU and shared test suite from MPS build
+        if: ${{ !cancelled() && steps.test_data.outcome == 'success' }}
+        run: |
+          build-mps/tests/ctranslate2_test tests/data \
+            --gtest_filter='-MPSBackendTest.*:MPSPrimitiveTest.*:TranslatorTest.MPS*:MPS/*' \
+            --gtest_output=xml:cpu-results.xml
+
+      - name: Build MPS Python bindings
+        id: build_python
+        if: ${{ !cancelled() && steps.test_data.outcome == 'success' }}
+        env:
+          CTRANSLATE2_ROOT: ${{ github.workspace }}/install-mps
+          CMAKE_BUILD_PARALLEL_LEVEL: 5
+        run: |
+          python -m pip install -r python/install_requirements.txt pytest
+          python -m pip install ./python
+
+      - name: Test Python CPU selection and explicit MPS inference
+        if: ${{ !cancelled() && steps.build_python.outcome == 'success' }}
+        env:
+          CT2_TEST_MPS: 1
+        run: |
+          python -m pytest --import-mode=importlib -v python/tests/test_mps.py \
+            --junitxml=python-mps-results.xml
+
+      - name: Upload hardware details and test reports
+        if: always()
+        uses: actions/upload-artifact@v4
+        with:
+          name: mps-macos-15-results
+          retention-days: 7
+          path: |
+            runner-hardware.txt
+            metal-probe.txt
+            mps-results.xml
+            cpu-results.xml
+            python-mps-results.xml
+
   build-python-wheels:
     runs-on: ${{ matrix.os }}
     strategy:

.gitignore

@@ -1,6 +1,10 @@
 *.pyc
 /.vs
 /build
+/build-*/
+/install-*/
+**/.DS_Store
+python/ctranslate2/_ext*.so
 
 CMake*.json
 .idea

CHANGELOG.md

@@ -1,5 +1,9 @@
 ## [Unreleased]
 
+### New features
+
+* Add an experimental Apple Metal (MPS) backend for Apple Silicon GPUs, selected with `device="mps"` (#2077) by [@TBO22](https://github.com/TBO22). It is opt-in at build time (`-DWITH_MPS=ON`), it is not included in the published wheels, and `device="auto"` never selects it
+
 ## [v4.8.3](https://github.com/OpenNMT/CTranslate2/releases/tag/v4.8.3) (2026-10-13)
 
 ### New features

CMakeLists.txt

@@ -15,6 +15,7 @@ option(WITH_RUY "Compile with Ruy backend" OFF)
 option(WITH_CUDA "Compile with CUDA backend" OFF)
 option(WITH_CUDNN "Compile with cuDNN backend" OFF)
 option(WITH_HIP "Compile with AMD HIP GPU backend" OFF)
+option(WITH_MPS "Compile with Apple Metal Performance Shaders backend" OFF)
 option(CUDA_DYNAMIC_LOADING "Dynamically load CUDA libraries at runtime" OFF)
 option(ENABLE_CPU_DISPATCH "Compile CPU kernels for multiple ISA and dispatch at runtime" ON)
 option(ENABLE_PROFILING "Compile with profiling support" OFF)
@@ -65,7 +66,13 @@ endif(NOT CMAKE_BUILD_TYPE)
 set(CMAKE_CXX_STANDARD 17)
 
 if(APPLE)
-  set(CMAKE_OSX_DEPLOYMENT_TARGET 10.13)
+  if(WITH_MPS)
+    # Apple Silicon Macs ship with macOS 11 or newer. This also covers the
+    # Metal APIs used by the persistent command stream and MPSMatrix paths.
+    set(CMAKE_OSX_DEPLOYMENT_TARGET 11.0)
+  else()
+    set(CMAKE_OSX_DEPLOYMENT_TARGET 10.13)
+  endif()
 endif()
 
 
@@ -786,6 +793,44 @@ else()
   add_library(${PROJECT_NAME} ${SOURCES})
 endif()
 
+if(WITH_MPS)
+  if(NOT APPLE)
+    message(FATAL_ERROR "WITH_MPS=ON is only supported on Apple platforms")
+  endif()
+  if(WITH_CUDA OR WITH_HIP)
+    message(FATAL_ERROR "WITH_MPS cannot be enabled together with CUDA or HIP")
+  endif()
+
+  message(STATUS "Enabling MPS (Metal) backend for Apple Silicon")
+  enable_language(OBJCXX)
+  target_compile_definitions(${PROJECT_NAME} PUBLIC CT2_WITH_MPS)
+
+  set(MPS_SOURCES
+    src/mps/allocator.mm
+    src/mps/kernels.mm
+    src/mps/ops_impl.mm
+    src/mps/ops_stubs.mm
+    src/mps/primitives.mm
+    src/mps/utils.mm
+    src/ops/conv1d_mps.mm
+    src/ops/layer_norm_mps.mm
+    src/ops/mean_mps.mm
+    src/ops/rms_norm_mps.mm
+    src/ops/rotary_mps.mm
+    src/ops/softmax_mps.cc
+  )
+  target_sources(${PROJECT_NAME} PRIVATE ${MPS_SOURCES})
+
+  find_library(METAL_LIBRARY Metal REQUIRED)
+  find_library(MPS_LIBRARY MetalPerformanceShaders REQUIRED)
+  find_library(FOUNDATION_LIBRARY Foundation REQUIRED)
+  target_link_libraries(${PROJECT_NAME} PRIVATE
+    ${METAL_LIBRARY}
+    ${MPS_LIBRARY}
+    ${FOUNDATION_LIBRARY}
+  )
+endif()
+
 include(GenerateExportHeader)
 generate_export_header(${PROJECT_NAME})
 set_property(TARGET ${PROJECT_NAME} PROPERTY VERSION ${CTRANSLATE2_VERSION})

README.md

@@ -60,6 +60,8 @@ See the [documentation](https://opennmt.net/CTranslate2) for more information an
 
 If you have an AMD ROCm GPU, we provide specific Python wheels on the [releases page](https://github.com/OpenNMT/CTranslate2/releases/).
 
+**Apple Metal (experimental).** Apple Silicon GPUs can be used with `device="mps"`. This backend is opt-in: it is only available when building from source with `-DWITH_MPS=ON`, it is not included in the published wheels, and `device="auto"` never selects it. See [hardware support](docs/hardware_support.md) for the supported features.
+
 ## Web Server
 
 [ctranslate2-web-server](https://github.com/jordimas/ctranslate2-web-server) is a web server built on top of CTranslate2 that exposes an OpenAI-compatible REST API, making it easy to integrate CTranslate2 models into applications that already support the OpenAI API.

cli/translator.cc

@@ -30,7 +30,7 @@ int main(int argc, char* argv[]) {
      cxxopts::value<size_t>()->default_value("1"))
     ("intra_threads", "Number of computation threads (set to 0 to use the default value).",
      cxxopts::value<size_t>()->default_value("0"))
-    ("device", "Device to use (can be cpu, cuda, auto).",
+    ("device", "Device to use (can be cpu, cuda, mps (experimental), auto).",
      cxxopts::value<std::string>()->default_value("cpu"))
     ("device_index", "Comma-separated list of device IDs to use.",
      cxxopts::value<std::vector<int>>()->default_value("0"))
@@ -44,6 +44,8 @@ int main(int argc, char* argv[]) {
      cxxopts::value<std::string>()->default_value("default"))
     ("cuda_compute_type", "Computation type on CUDA devices (overrides compute_type)",
      cxxopts::value<std::string>())
+    ("mps_compute_type", "Computation type on MPS devices (overrides compute_type)",
+     cxxopts::value<std::string>())
     ("cpu_compute_type", "Computation type on CPU devices (overrides compute_type)",
      cxxopts::value<std::string>())
     ;
@@ -139,6 +141,10 @@ int main(int argc, char* argv[]) {
     if (args.count("cuda_compute_type"))
       compute_type = ctranslate2::str_to_compute_type(args["cuda_compute_type"].as<std::string>());
     break;
+  case ctranslate2::Device::MPS:
+    if (args.count("mps_compute_type"))
+      compute_type = ctranslate2::str_to_compute_type(args["mps_compute_type"].as<std::string>());
+    break;
   };
 
   ctranslate2::ReplicaPoolConfig pool_config;

docs/environment_variables.md

@@ -51,6 +51,71 @@ export CT2_CUDA_CACHING_ALLOCATOR_CONFIG=8,3,7,6291455
 
 See the description of each parameter in the [allocator implementation](https://github.com/NVIDIA/cub/blob/main/cub/util_allocator.cuh).
 
+```{attention}
+**Apple Metal (experimental).** The `CT2_MPS_*` variables below configure the experimental Apple Metal (MPS) backend. This backend is opt-in at build time (`-DWITH_MPS=ON`), it is not included in the published Python wheels, and `device="auto"` never selects it. The variables and their defaults may change.
+```
+
+## `CT2_MPS_MAX_COMPUTE_OPS`
+
+Set the maximum number of compute operations encoded before the current MPS
+command buffer is committed without waiting. The default is `16`.
+
+## `CT2_MPS_MAX_BLIT_OPS`
+
+Set the maximum number of copy/blit operations encoded before the current MPS
+command buffer is committed without waiting. The default is `64`.
+
+## `CT2_MPS_MAX_INFLIGHT_COMMAND_BUFFERS`
+
+Set the maximum number of unfinished MPS command buffers. The default is `8`,
+which keeps useful CPU/GPU overlap while bounding Metal temporary resources on
+unified-memory systems. Increasing this value can help large-memory devices,
+but should be validated with cold, end-to-end inference rather than GEMM-only
+benchmarks.
+
+## `CT2_MPS_MAX_OPS`
+
+Legacy command-buffer operation limit. When one of the more specific MPS
+compute or blit limits above is not set, this value overrides its default. It
+does not change `CT2_MPS_MAX_INFLIGHT_COMMAND_BUFFERS`.
+
+## `CT2_MPS_PROFILE`
+
+Print aggregated MPS execution counters at process exit, including command
+buffers, dispatches, synchronization points, GEMM paths, TopK calls, copies,
+allocations, and buffer lookups. Profiling is disabled by default.
+
+## `CT2_MPS_LOG_GEMM`
+
+Log each MPS GEMM shape, transpose and stride configuration, and selected
+execution path. This option is intended for debugging and benchmarking.
+
+## `CT2_MPS_LOG_SYNC`
+
+Log each MPS synchronization together with its reason and host-visible byte
+count.
+
+## `CT2_MPS_USE_GEMV`
+
+Enable the specialized MPS batch-size-1 GEMV path (default). Set to `0` to
+force the general GEMM path for correctness comparisons or benchmarking.
+
+## `CT2_MPS_USE_TOPK`
+
+Enable the MPS GPU TopK path (default). Set to `0` to use the fallback path for
+debugging or performance comparisons.
+
+## `CT2_MPS_CACHE_INT8_FP16`
+
+Resolve `int8` and `int8_float16` requests on MPS to the faster and more stable
+FP16 execution path (default). Stored INT8 weights are expanded once while the
+model is loaded; the compressed device allocation is not retained.
+
+Set this variable to `0` to explicitly keep weights compressed and use the
+experimental native MPS INT8 path. This reduces model memory, but can be slower
+than FP16 and may produce different decoding decisions because activations are
+quantized at runtime.
+
 ## `CT2_FORCE_CPU_ISA`
 
 Force CTranslate2 to select a specific instruction set architecture (ISA). Possible values are:

docs/hardware_support.md

@@ -17,6 +17,26 @@ See the [environment variables](environment_variables.md) `CT2_USE_MKL` and `CT2
 
 ## GPU
 
+### NVIDIA CUDA
+
 * NVIDIA GPUs with a Compute Capability greater or equal to 3.5
 
 The driver requirement depends on the CUDA version. See the [CUDA Compatibility guide](https://docs.nvidia.com/deploy/cuda-compatibility/index.html) for more information.
+
+### Apple Metal Performance Shaders
+
+The experimental MPS backend supports Apple Silicon Macs running macOS 11 or
+newer. It is not included in the prebuilt Python wheels: build from source
+with `-DWITH_MPS=ON`, then select `device="mps"` in the C++ or Python API.
+
+The backend supports FP32, FP16, BF16, and experimental signed INT8
+computation. FP16 is the recommended compute type and is selected by
+`compute_type="auto"`. Requests for `int8` and `int8_float16` also resolve to
+FP16 by default because native INT8 is slower on current Apple GPUs and can
+change decoding decisions. See [Quantization](quantization.md) for the native
+INT8 opt-in and exact type conversions, and [Installation](installation.md)
+for a build example.
+
+The MPS backend cannot be enabled together with CUDA or HIP in the same build.
+FlashAttention, AWQ INT4, INT16 GEMM, distributed collectives, and
+packed/shifted-u8 INT8 GEMM are not currently implemented on MPS.