refactor: split operator headers into public and detail layers

feat: add MACA-based MetaX backend support
Co-authored-by: wawahejun <hejunlbbc@gmail.com>
2026-05-08 11:18:47 +08:00 · 2026-05-08 09:49:13 +08:00 · 2026-05-05 16:50:18 +08:00 · 2026-05-05 16:50:18 +08:00 · 2026-05-05 16:50:18 +08:00 · 2026-05-05 16:50:18 +08:00
82 changed files with 1580 additions and 633 deletions
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@ -10,14 +10,37 @@ set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 set(CMAKE_POSITION_INDEPENDENT_CODE ON)

-list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+list(PREPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")

-if(CAMP_ENABLE_NVIDIA)
-  if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-    set(CMAKE_CUDA_ARCHITECTURES native)
+if(CAMP_ENABLE_NVIDIA AND CAMP_ENABLE_METAX)
+  message(FATAL_ERROR "CAMP_ENABLE_NVIDIA and CAMP_ENABLE_METAX cannot both be ON; build separate variants")
+endif()
+
+if(CAMP_ENABLE_METAX)
+  find_package(MetaX REQUIRED)
+  set(ENV{CUCC_PATH} "${MetaX_CUBRIDGE_ROOT}")
+  set(ENV{CUDA_PATH} "${MetaX_CUBRIDGE_ROOT}")
+  list(APPEND CMAKE_MODULE_PATH "${MetaX_CMAKE_MODULE_DIR}")
+  set(CMAKE_MACA_SOURCE_FILE_EXTENSIONS maca;cu;cc;cpp;c)
+endif()
+
+if(CAMP_ENABLE_NVIDIA OR CAMP_ENABLE_METAX)
+  if(CAMP_ENABLE_METAX)
+    enable_language(MACA)
+  else()
+    if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+      set(CMAKE_CUDA_ARCHITECTURES native)
+    endif()
+    enable_language(CUDA)
+    include(cmake/cuda_helpers.cmake)
  endif()
-  enable_language(CUDA)
-  include(cmake/cuda_helpers.cmake)
+endif()
+
+if(CAMP_ENABLE_METAX)
+  set(CAMP_METAX_COMPILE_OPTIONS
+    -D__MACA_NO_HALF_OPERATORS__
+    -D__FAST_BLOCK_RED__
+  )
 endif()

 add_subdirectory(ops)
--- a/README.md
+++ b/README.md
@ -9,15 +9,29 @@ small, but its workflow mirrors production operator libraries:
 4. Expose a Python API with out-of-place, out-variant, and prepared execution.
 5. Validate correctness against PyTorch and benchmark steady-state execution.

+## Python Layout
+
+```text
+python/
+  operator_runtime/
+    backend.py
+    ops/
+    _internal/
+  operator_runtime_testing/
+```
+
+- `operator_runtime.ops` contains public operator bindings.
+- `operator_runtime._internal` contains private FFI/runtime plumbing.
+- `operator_runtime_testing` contains test-only helpers such as assertions and benchmark utilities.

 ## Operators

 | Operator | NVIDIA C++ | TileLang | MetaX |
 | --- | --- | --- | --- |
-| `copy` | runnable | runnable when TileLang is installed | stub |
-| `vector_add` | runnable | runnable when TileLang is installed | stub |
-| `reduce_sum` | runnable, row-wise fp32 | runnable when TileLang is installed | stub |
-| `softmax` | runnable, row-wise fp32 | runnable when TileLang is installed | stub |
+| `copy` | runnable | runnable when TileLang is installed | runnable |
+| `vector_add` | runnable | runnable when TileLang is installed | runnable |
+| `reduce_sum` | runnable, row-wise fp32 | runnable when TileLang is installed | runnable, row-wise fp32 |
+| `softmax` | runnable, row-wise fp32 | runnable when TileLang is installed | runnable, row-wise fp32 |

 ## Setup

@ -34,6 +48,10 @@ cmake .. -DCAMP_ENABLE_NVIDIA=ON -DCAMP_ENABLE_METAX=OFF
 cmake --build . -j$(nproc)
 ```

+```bash
+./scripts/build_metax.sh build
+```
+
 ## Validate

 ```bash
@ -41,28 +59,17 @@ python tests/run_ops.py --op copy --backend nvidia --mode all
 CAMP_BUILD_DIR=build pytest tests/ -v --backend nvidia
 pytest tests/ -v --backend tilelang
 python tests/run_ops.py --op all --backend nvidia --mode bench
+./scripts/build_metax.sh test
 ```

-The TileLang backend requires the `tilelang` Python package.
-
-## Adding a New Operator
-
-1. Create `ops/<name>/nvidia/<name>_cuda.h` with the C API (4 functions: create, workspace, execute, destroy).
-2. Create `ops/<name>/nvidia/<name>_cuda.cu` with the CUDA implementation.
-3. Create `python/operator_runtime/ops/<name>.py` using `ctypes_bindings.bind_*` functions.
-4. Create `tests/cases/<name>.py` with `correctness_cases()`, `api_error_cases()`, and `benchmark_cases()`.
-5. Create `tests/ops/test_<name>.py` and `tests/bench/<name>.py`.
-6. Re-run `cmake ..` in the build directory (the glob will pick up the new `.cu` file).
-7. Register the public API in `python/operator_runtime/__init__.py`.
-
-No YAML, no code generation, no registration step.
+The TileLang backend requires the `tilelang` Python package. The MetaX backend is built as a separate variant, should use `CAMP_ENABLE_NVIDIA=OFF`, and stores backend sources under `ops/*/metax/*.maca`.

 ## Production Mapping

 | Training concept | Production equivalent |
 | --- | --- |
 | directory convention `ops/<op>/nvidia/*.cu` | build system auto-discovery / operator registry |
-| C header `ops/<op>/nvidia/<op>_cuda.h` | reviewed operator API contract |
+| C header `include/operator_runtime/ops/<op>.h` | reviewed operator API contract |
 | descriptor lifecycle | create, workspace, execute, destroy |
 | `tests/cases/<op>.py` | correctness, layout, and API contract coverage |
 | `PerformanceResult` | profiler report row with latency, bytes, flops, bandwidth |
--- a/README.zh.md
+++ b/README.zh.md
@ -8,14 +8,29 @@
 4. 提供 Python API，支持 out-of-place、out-variant 和 prepared 执行。
 5. 通过 PyTorch 做正确性验证，并做稳定态性能测试。

+## Python 目录结构
+
+```text
+python/
+  operator_runtime/
+    backend.py
+    ops/
+    _internal/
+  operator_runtime_testing/
+```
+
+- `operator_runtime.ops` 放公开算子绑定。
+- `operator_runtime._internal` 放私有 FFI 和运行时细节。
+- `operator_runtime_testing` 放断言、benchmark 等仅测试使用的工具。
+
 ## 算子

 | 算子 | NVIDIA C++ | TileLang | MetaX |
 | --- | --- | --- | --- |
-| `copy` | 可运行 | 安装 TileLang 后可运行 | stub |
-| `vector_add` | 可运行 | 安装 TileLang 后可运行 | stub |
-| `reduce_sum` | 可运行，row-wise fp32 | 安装 TileLang 后可运行 | stub |
-| `softmax` | 可运行，row-wise fp32 | 安装 TileLang 后可运行 | stub |
+| `copy` | 可运行 | 安装 TileLang 后可运行 | 可运行 |
+| `vector_add` | 可运行 | 安装 TileLang 后可运行 | 可运行 |
+| `reduce_sum` | 可运行，row-wise fp32 | 安装 TileLang 后可运行 | 可运行，row-wise fp32 |
+| `softmax` | 可运行，row-wise fp32 | 安装 TileLang 后可运行 | 可运行，row-wise fp32 |

 ## 安装

@ -32,6 +47,10 @@ cmake .. -DCAMP_ENABLE_NVIDIA=ON -DCAMP_ENABLE_METAX=OFF
 cmake --build . -j$(nproc)
 ```

+```bash
+./scripts/build_metax.sh build
+```
+
 ## 验证

 ```bash
@ -39,16 +58,17 @@ python tests/run_ops.py --op copy --backend nvidia --mode all
 CAMP_BUILD_DIR=build pytest tests/ -v --backend nvidia
 pytest tests/ -v --backend tilelang
 python tests/run_ops.py --op all --backend nvidia --mode bench
+./scripts/build_metax.sh test
 ```

-TileLang 后端需要安装 `tilelang` Python 包。
+TileLang 后端需要安装 `tilelang` Python 包。MetaX 后端使用独立构建产物，构建时应关闭 NVIDIA 变体，并将后端源码放在 `ops/*/metax/*.maca`。

 ## 生产映射

 | 训练概念 | 生产等价物 |
 | --- | --- |
 | `ops/<op>/nvidia/*.cu` 目录约定 | 构建系统自动发现 / 算子注册 |
-| `ops/<op>/nvidia/<op>_cuda.h` 头文件 | 经过评审的算子 API 契约 |
+| `include/operator_runtime/ops/<op>.h` 头文件 | 经过评审的算子 API 契约 |
 | descriptor 生命周期 | create、workspace、execute、destroy |
 | `tests/cases/<op>.py` | 正确性、布局和 API 契约覆盖 |
 | `PerformanceResult` | 包含延迟、字节数、FLOPs、带宽的 profiler 报表行 |
--- a/cmake/CMakeMACAInformation.cmake
+++ b/cmake/CMakeMACAInformation.cmake
@ -0,0 +1,16 @@
+set(CMAKE_MACA_COMPILE_OBJECT
+  "<CMAKE_MACA_COMPILER> <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> -c <SOURCE>"
+)
+
+set(CMAKE_INCLUDE_FLAG_MACA "-I")
+set(CMAKE_MACA_COMPILE_OPTIONS_PIC "-fPIC")
+
+set(CMAKE_MACA_CREATE_SHARED_LIBRARY
+  "<CMAKE_MACA_COMPILER> <LANGUAGE_COMPILE_FLAGS> <LINK_FLAGS> -shared -o <TARGET> <OBJECTS> <LINK_LIBRARIES> --maca-path=$ENV{MACA_PATH}"
+)
+
+set(CMAKE_MACA_LINK_EXECUTABLE
+  "<CMAKE_MACA_COMPILER> <FLAGS> <LINK_FLAGS> -o <TARGET> <OBJECTS> <LINK_LIBRARIES> --maca-path=$ENV{MACA_PATH}"
+)
+
+set(CMAKE_MACA_INFORMATION_LOADED 1)
--- a/cmake/FindMetaX.cmake
+++ b/cmake/FindMetaX.cmake
@ -1,6 +1,82 @@
-# Placeholder for a future MetaX toolchain integration.
-#
-# A production implementation should detect MACA_PATH, htcc/mxcc, runtime
-# libraries, include directories, and backend-specific source suffixes.
-set(MetaX_FOUND OFF)
+set(_MetaX_SEARCH_ROOTS "")

+if(DEFINED ENV{MACA_PATH})
+  list(APPEND _MetaX_SEARCH_ROOTS "$ENV{MACA_PATH}")
+endif()
+list(APPEND _MetaX_SEARCH_ROOTS /opt/maca)
+
+find_path(
+  MetaX_ROOT
+  NAMES include/mcr/mc_runtime.h
+  PATHS ${_MetaX_SEARCH_ROOTS}
+  NO_DEFAULT_PATH
+)
+
+find_program(
+  MetaX_MXCC_EXECUTABLE
+  NAMES mxcc
+  PATHS
+    ${MetaX_ROOT}/mxgpu_llvm/bin
+    /opt/maca/mxgpu_llvm/bin
+)
+
+find_program(
+  MetaX_CUCC_EXECUTABLE
+  NAMES cucc
+  PATHS
+    ${MetaX_ROOT}/tools/cu-bridge/bin
+    /opt/maca/tools/cu-bridge/bin
+)
+
+find_path(
+  MetaX_INCLUDE_DIR
+  NAMES mcr/mc_runtime.h
+  PATHS
+    ${MetaX_ROOT}/include
+    /opt/maca/include
+)
+
+find_library(
+  MetaX_MCRUNTIME_LIBRARY
+  NAMES mcruntime libmcruntime.so
+  PATHS
+    ${MetaX_ROOT}/lib
+    /opt/maca/lib
+)
+
+find_library(
+  MetaX_MXCRT_LIBRARY
+  NAMES mxc-runtime64 libmxc-runtime64.so
+  PATHS
+    ${MetaX_ROOT}/lib
+    /opt/maca/lib
+)
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(
+  MetaX
+  REQUIRED_VARS
+    MetaX_ROOT
+    MetaX_MXCC_EXECUTABLE
+    MetaX_INCLUDE_DIR
+    MetaX_MCRUNTIME_LIBRARY
+)
+
+if(MetaX_FOUND)
+  set(MetaX_CUBRIDGE_ROOT "${MetaX_ROOT}/tools/cu-bridge")
+  set(MetaX_CMAKE_MODULE_DIR "${MetaX_CUBRIDGE_ROOT}/cmake_module/maca")
+  set(MetaX_INCLUDE_DIRS "${MetaX_INCLUDE_DIR}")
+  set(MetaX_LIBRARIES "${MetaX_MCRUNTIME_LIBRARY}")
+  if(MetaX_MXCRT_LIBRARY)
+    list(APPEND MetaX_LIBRARIES "${MetaX_MXCRT_LIBRARY}")
+  endif()
+endif()
+
+mark_as_advanced(
+  MetaX_ROOT
+  MetaX_MXCC_EXECUTABLE
+  MetaX_CUCC_EXECUTABLE
+  MetaX_INCLUDE_DIR
+  MetaX_MCRUNTIME_LIBRARY
+  MetaX_MXCRT_LIBRARY
+)
--- a/docs/how-to-add-an-operator.md
+++ b/docs/how-to-add-an-operator.md
@ -1,7 +1,5 @@
 # 如何开发一个新算子

-本文只描述当前仓库里新增算子的实际流程，不展开代码细节。
-
 ## 目标

 在当前项目里，新增一个可运行算子的最小闭环包括四部分：
@ -11,15 +9,6 @@
 3. 在 `tests/` 下补正确性测试和 benchmark 入口。
 4. 重新构建并验证。

-## 当前仓库的真实约束
-
- 现在没有 `operator.yaml`。
- 现在没有代码生成步骤。
- 现在没有单独的 operator registry 文件。
- NVIDIA 的 `.cu` 文件通过目录约定自动发现。
- 新增或重命名 `.cu` 后，需要重新执行一次 `cmake ..`。
-
-也就是说，新增算子主要依赖目录结构和命名约定，而不是额外的注册配置。

 ## Step 1：明确算子接口

@ -43,7 +32,7 @@
 - `ops/<算子名>/nvidia/`
 - `python/operator_runtime/ops/<算子名>.py`
 - `tests/cases/<算子名>.py`
- `tests/ops/test_<算子名>.py`
+- `tests/op_tests/test_<算子名>.py`
 - `tests/bench/<算子名>.py`

 如果后续要支持 TileLang 或 MetaX，再分别补 `tilelang/` 或 `metax/`。
@ -123,7 +112,7 @@ Python 入口放在 `python/operator_runtime/ops/<算子名>.py`。

 ## Step 7：补正确性测试

-`tests/ops/test_<算子名>.py` 主要负责三件事：
+`tests/op_tests/test_<算子名>.py` 主要负责三件事：

 1. 正确性对比
 2. API contract 检查
@ -202,7 +191,7 @@ API contract 测试主要覆盖 shape 不匹配、dtype 不匹配、非 contiguo
 2. `python/operator_runtime/ops/<算子名>.py` 已补齐。
 3. 两个 `__init__.py` 已导出新接口。
 4. `tests/cases/<算子名>.py` 已补数据。
-5. `tests/ops/test_<算子名>.py` 已补测试。
+5. `tests/op_tests/test_<算子名>.py` 已补测试。
 6. `tests/bench/<算子名>.py` 已补 benchmark。
 7. 新增 `.cu` 后已经重新执行过 `cmake ..`。
 8. 至少完成一次单算子验证。
--- a/docs/phase1-kernel-writing.md
+++ b/docs/phase1-kernel-writing.md
@ -0,0 +1,103 @@
+# 第一阶段训练目标：写 Kernel
+
+## 目标
+
+本阶段主要需要做一件事：
+
+**给现有算子写 kernel 文件。**
+
+其余所有代码（descriptor 生命周期、Python 绑定、测试、构建）都已经写好，剩下需要补全计算逻辑本身。
+
+默认要求是：
+
+1. **优先把 kernel 写对、跑通。**
+2. 不需要一开始就追求最优性能。
+3. 先完成 `kernel.cuh` / `kernel.py` 里的 TODO，再考虑更激进的优化。
+
+如果你已经完成基础目标，并且学有余力，希望进一步冲性能，那么可以继续修改对应算子的包装层文件（例如 NVIDIA 后端的 `.cu` 文件），调整 launch policy、切换不同 kernel、加入更激进的 specialization。这些都属于进阶优化内容，不属于第一阶段的必做要求。
+
+---
+
+## 需要写的文件
+
+每个算子有两种后端，对应两个 kernel 文件：
+
+| 算子 | NVIDIA kernel | TileLang kernel |
+|------|--------------|-----------------|
+| `copy` | `ops/copy/nvidia/kernel.cuh` | `ops/copy/tilelang/kernel.py` |
+| `vector_add` | `ops/vector_add/nvidia/kernel.cuh` | `ops/vector_add/tilelang/kernel.py` |
+| `reduce_sum` | `ops/reduce_sum/nvidia/kernel.cuh` | `ops/reduce_sum/tilelang/kernel.py` |
+| `softmax` | `ops/softmax/nvidia/kernel.cuh` | `ops/softmax/tilelang/kernel.py` |
+
+建议按从简到难的顺序：copy → vector_add → reduce_sum → softmax。
+
+---
+
+## NVIDIA Kernel（`.cuh` 文件）
+
+### 写什么
+
+一个 `__global__` 函数，放在对应算子的命名空间下。函数只负责计算逻辑，输入输出是裸指针，不涉及任何 descriptor 或 API。
+
+本阶段默认**不需要修改**同目录下的 `.cu` 文件。仓库已经提供了兼容性优先的 launch 样板，这一阶段只要求把 `kernel.cuh` 写对、跑通。
+
+如果你已经完成基础目标，并且想继续冲性能，可以把 `.cu` 当作进阶优化层：在那里调整线程数、切换不同 kernel、加入更激进的 specialization。但这些不属于第一阶段的必做内容。
+
+### 四个算子的难度递进
+
+**copy / vector_add**：逐元素操作，用 grid-stride loop 模式。每个线程负责若干个独立的元素，线程之间没有通信。
+
+**reduce_sum**：行规约。每行一个 block，block 内线程先各自累加自己负责的列，再通过 shared memory 做树形归约，最终 `smem[0]` 是整行的结果。需要用 `__syncthreads()` 同步。
+
+**softmax**：基础版本使用三段流程：第一遍求行最大值（数值稳定性），第二遍求 `exp(x - max)` 的和并把中间结果写到输出，第三遍再除以 sum。每一遍都需要 block 内同步。
+
+### 需要理解的概念
+
+- **grid-stride loop**：为什么这样写可以处理任意大小的 tensor
+- **shared memory 规约**：树形归约的每一步在做什么，为什么需要 `__syncthreads()`
+- **softmax 减 max**：为什么直接算 `exp(x)` 会出问题，减去行最大值为什么不改变结果
+
+---
+
+## TileLang Kernel（`.py` 文件）
+
+### 写什么
+
+一个用 `@tilelang.jit` 装饰的 Python 函数，用 TileLang DSL 描述 tile 粒度的计算逻辑。TileLang 会把它编译成真正的 CUDA kernel。
+
+和 NVIDIA 一样，本阶段默认只需要填写 `kernel.py` 里的 TODO，不需要修改外层适配代码。
+
+### 四个算子的难度递进
+
+**copy**：直接用内置 `T.copy` 在两个 tile 之间搬数据，无需手写循环。
+
+**vector_add**：用 `T.Parallel` 循环在 tile 内逐元素计算，理解 `T.Parallel` 与 `T.Serial` 的区别。
+
+**reduce_sum**：需要分块累加。外层用 `T.Serial` 顺序遍历列方向的分块（因为要累积状态），内层用 `T.reduce_sum` 对 fragment 做规约。
+
+**softmax**：两遍扫描，使用 online softmax 算法。第一遍滚动维护 log-sum-exp 状态，第二遍用最终的 lse 归一化。计算中用 `exp2` / `log2` 替代 `exp` / `log`。
+
+### 需要理解的概念
+
+- **`T.Parallel` vs `T.Serial`**：什么情况下循环内的迭代可以并行，什么情况下必须顺序
+- **`T.alloc_fragment`**：tile 级别的局部 buffer，对应寄存器或 shared memory
+- **`T.copy`**：把全局内存的一块搬到 fragment，不是逐元素赋值
+- **online softmax**：为什么一遍扫描就能算出正确的归一化，log-sum-exp 的滚动更新逻辑
+- **`exp2` / `log2`**：为什么 TileLang 里用这两个而不是 `exp` / `log`
+
+---
+
+## 验证方式
+
+每写完一个 kernel，用对应的测试验证：
+
+```bash
+# NVIDIA kernel（需要先重新编译）
+cd build && cmake .. && cmake --build . -j$(nproc)
+CAMP_BUILD_DIR=/workspace/build pytest tests/op_tests/test_<算子名>.py -v --backend nvidia
+
+# TileLang kernel
+CAMP_BUILD_DIR=/workspace/build pytest tests/op_tests/test_<算子名>.py -v --backend tilelang
+```
+
+四个算子两种后端全部通过，阶段一完成。
--- a/examples/01_copy.py
+++ b/examples/01_copy.py
@ -15,7 +15,7 @@ from operator_runtime import copy

 def main() -> None:
    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang"])
+    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
    args = parser.parse_args()

    src = torch.randn((1024,), device="cuda", dtype=torch.float32)
--- a/examples/02_vector_add.py
+++ b/examples/02_vector_add.py
@ -15,7 +15,7 @@ from operator_runtime import vector_add

 def main() -> None:
    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang"])
+    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
    args = parser.parse_args()

    a = torch.randn((1024,), device="cuda", dtype=torch.float32)
--- a/examples/03_reduce_sum.py
+++ b/examples/03_reduce_sum.py
@ -15,7 +15,7 @@ from operator_runtime import reduce_sum

 def main() -> None:
    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang"])
+    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
    args = parser.parse_args()

    src = torch.randn((32, 128), device="cuda", dtype=torch.float32)
--- a/examples/04_softmax.py
+++ b/examples/04_softmax.py
@ -15,7 +15,7 @@ from operator_runtime import softmax

 def main() -> None:
    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang"])
+    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
    args = parser.parse_args()

    src = torch.randn((32, 128), device="cuda", dtype=torch.float32)
--- a/include/operator_runtime/detail/cuda_helpers.h
+++ b/include/operator_runtime/detail/cuda_helpers.h
@ -5,12 +5,12 @@
 #ifdef __CUDACC__
 #include <cuda_runtime.h>

-#define OPRT_CUDA_RETURN_IF_ERROR(expr)        \
-    do {                                       \
-        cudaError_t err__ = (expr);            \
-        if (err__ != cudaSuccess) {            \
-            return OPRT_ERR_RUNTIME;           \
-        }                                      \
+#define OPRT_CUDA_RETURN_IF_ERROR(expr) \
+    do { \
+        cudaError_t err__ = (expr); \
+        if (err__ != cudaSuccess) { \
+            return OPRT_ERR_RUNTIME; \
+        } \
    } while (0)

 namespace oprt {
@ -27,4 +27,3 @@ inline int blocks_for(int64_t n, int threads) {
 } // namespace oprt

 #endif
-
--- a/include/operator_runtime/detail/elementwise.h
+++ b/include/operator_runtime/detail/elementwise.h
@ -21,4 +21,3 @@ inline bool elementwise_fast_path(const oprt_tensor_view_t &out,
 } // namespace oprt

 #endif
-
--- a/include/operator_runtime/detail/operation.h
+++ b/include/operator_runtime/detail/operation.h
@ -1,7 +1,5 @@
 #pragma once

-#ifdef __cplusplus
-
 #include <string>

 namespace oprt {
@ -13,6 +11,3 @@ struct OperationSpec {
 };

 } // namespace oprt
-
-#endif
-
--- a/include/operator_runtime/detail/tensor_checks.h
+++ b/include/operator_runtime/detail/tensor_checks.h
@ -37,4 +37,3 @@ inline oprt_status_t check_same_shape(const oprt_tensor_view_t &a,
 } // namespace oprt

 #endif
-
--- a/include/operator_runtime/operator_runtime.h
+++ b/include/operator_runtime/operator_runtime.h
@ -0,0 +1,9 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/tensor_view.h"
+#include "operator_runtime/ops/copy.h"
+#include "operator_runtime/ops/reduce_sum.h"
+#include "operator_runtime/ops/softmax.h"
+#include "operator_runtime/ops/vector_add.h"
--- a/include/operator_runtime/ops/copy.h
+++ b/include/operator_runtime/ops/copy.h
@ -6,16 +6,16 @@
 extern "C" {
 #endif

-OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *dst,
    const oprt_tensor_view_t *src);

-OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia(
+OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size);

-OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia(
+OPRT_EXPORT oprt_status_t oprt_execute_copy(
    oprt_operator_descriptor_t desc,
    void *workspace,
    size_t workspace_size,
@ -23,10 +23,9 @@ OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia(
    const void *src,
    oprt_stream_t stream);

-OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor(
    oprt_operator_descriptor_t desc);

 #ifdef __cplusplus
 }
 #endif
-
--- a/include/operator_runtime/ops/reduce_sum.h
+++ b/include/operator_runtime/ops/reduce_sum.h
@ -6,17 +6,17 @@
 extern "C" {
 #endif

-OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *out,
    const oprt_tensor_view_t *in,
    int64_t axis);

-OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia(
+OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size);

-OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia(
+OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum(
    oprt_operator_descriptor_t desc,
    void *workspace,
    size_t workspace_size,
@ -24,10 +24,9 @@ OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia(
    const void *in,
    oprt_stream_t stream);

-OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor(
    oprt_operator_descriptor_t desc);

 #ifdef __cplusplus
 }
 #endif
-
--- a/ops/reduce_sum/nvidia/reduce_sum_cuda.h
+++ b/ops/reduce_sum/nvidia/reduce_sum_cuda.h
@ -6,17 +6,17 @@
 extern "C" {
 #endif

-OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *out,
    const oprt_tensor_view_t *in,
    int64_t axis);

-OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia(
+OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size);

-OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia(
+OPRT_EXPORT oprt_status_t oprt_execute_softmax(
    oprt_operator_descriptor_t desc,
    void *workspace,
    size_t workspace_size,
@ -24,10 +24,9 @@ OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia(
    const void *in,
    oprt_stream_t stream);

-OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor(
    oprt_operator_descriptor_t desc);

 #ifdef __cplusplus
 }
 #endif
-
--- a/include/operator_runtime/ops/vector_add.h
+++ b/include/operator_runtime/ops/vector_add.h
@ -6,17 +6,17 @@
 extern "C" {
 #endif

-OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *out,
    const oprt_tensor_view_t *a,
    const oprt_tensor_view_t *b);

-OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_nvidia(
+OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size);

-OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia(
+OPRT_EXPORT oprt_status_t oprt_execute_vector_add(
    oprt_operator_descriptor_t desc,
    void *workspace,
    size_t workspace_size,
@ -25,10 +25,9 @@ OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia(
    const void *b,
    oprt_stream_t stream);

-OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor_nvidia(
+OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor(
    oprt_operator_descriptor_t desc);

 #ifdef __cplusplus
 }
 #endif
-
--- a/ops/CMakeLists.txt
+++ b/ops/CMakeLists.txt
@ -1,5 +1,6 @@
 # Convention: operator CUDA sources auto-discovered by directory layout:
 #   ops/<operator>/nvidia/*.cu  ->  compiled when CAMP_ENABLE_NVIDIA is ON
+#   ops/<operator>/metax/*.maca ->  compiled when CAMP_ENABLE_METAX is ON
 # After adding a new operator, re-run cmake to pick up new files.

 set(CAMP_COMMON_SOURCES
@ -13,14 +14,31 @@ if(CAMP_ENABLE_NVIDIA)
  )
 endif()

-set(CAMP_OPERATOR_SOURCES ${CAMP_COMMON_SOURCES} ${CAMP_NVIDIA_SOURCES})
+set(CAMP_METAX_SOURCES "")
+if(CAMP_ENABLE_METAX)
+  file(GLOB_RECURSE CAMP_METAX_SOURCES
+    ${CMAKE_CURRENT_SOURCE_DIR}/*/metax/*.maca
+  )
+  set_source_files_properties(${CAMP_METAX_SOURCES} PROPERTIES LANGUAGE MACA)
+endif()
+
+set(CAMP_OPERATOR_SOURCES ${CAMP_COMMON_SOURCES} ${CAMP_NVIDIA_SOURCES} ${CAMP_METAX_SOURCES})

 add_library(camp_ops SHARED ${CAMP_OPERATOR_SOURCES})
-target_include_directories(camp_ops PUBLIC
-  "${CMAKE_CURRENT_SOURCE_DIR}/.."
-  "${CMAKE_CURRENT_SOURCE_DIR}/../include"
+target_include_directories(camp_ops
+  PUBLIC
+    "${CMAKE_CURRENT_SOURCE_DIR}/../include"
+  PRIVATE
+    "${CMAKE_CURRENT_SOURCE_DIR}/.."
 )

+if(CAMP_ENABLE_METAX)
+  target_include_directories(camp_ops PRIVATE ${MetaX_INCLUDE_DIRS})
+  target_link_libraries(camp_ops PRIVATE ${MetaX_LIBRARIES})
+  target_compile_options(camp_ops PRIVATE ${CAMP_METAX_COMPILE_OPTIONS})
+  set_target_properties(camp_ops PROPERTIES LINKER_LANGUAGE MACA)
+endif()
+
 target_compile_definitions(camp_ops PRIVATE
  $<$<BOOL:${CAMP_ENABLE_NVIDIA}>:CAMP_ENABLE_NVIDIA=1>
  $<$<BOOL:${CAMP_ENABLE_METAX}>:CAMP_ENABLE_METAX=1>
--- a/ops/common/metax/README.md
+++ b/ops/common/metax/README.md
@ -1,13 +1,12 @@
-# MetaX Backend Stub
+# MetaX Backend

-The training runtime keeps MetaX interfaces aligned with NVIDIA interfaces, but
-does not build MetaX code by default.
+The training runtime keeps MetaX interfaces aligned with the compiled C ABI
+backends and builds MetaX as a separate variant through MACA/cu-bridge.

-A production integration should add:
-
- MACA SDK discovery (`MACA_PATH`)
- htcc/mxcc compiler rules
- `.maca` source handling
- hcruntime/mcruntime link rules
- backend-specific CI on MetaX hardware
+Current expectations:

+- discover the MACA SDK from `MACA_PATH` or `/opt/maca`
+- build with `cmake_maca` and `ninja_maca`
+- compile `ops/*/metax/*.maca` into `libcamp_ops.so`
+- load `_metax` ABI symbols from Python FFI
+- validate on a MetaX-enabled PyTorch environment that exposes `cuda:0`
--- a/ops/copy/metax/README.md
+++ b/ops/copy/metax/README.md
@ -1,5 +1,4 @@
-# Copy MetaX Stub
-
-The copy ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and htcc/mxcc
-rules when MetaX hardware is available.
+# Copy MetaX Backend

+The copy MetaX backend exports the standard descriptor lifecycle with `_metax`
+symbols and is built from `ops/copy/metax/copy_metax.maca`.
--- a/ops/copy/metax/copy_metax.h
+++ b/ops/copy/metax/copy_metax.h
@ -1,6 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
-
--- a/ops/copy/metax/copy_metax.maca
+++ b/ops/copy/metax/copy_metax.maca
@ -0,0 +1,119 @@
+#include "operator_runtime/ops/copy.h"
+
+#include "operator_runtime/detail/cuda_helpers.h"
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/detail/elementwise.h"
+#include "operator_runtime/detail/tensor_checks.h"
+
+#include <cuda_fp16.h>
+
+namespace {
+
+struct CopyDescriptor final : oprt_operator_descriptor {
+    oprt_tensor_view_t dst_view;
+    oprt_tensor_view_t src_view;
+    int64_t elements = 0;
+
+    const char *op_name() const override {
+        return "copy";
+    }
+};
+
+template <typename T>
+__global__ void copy_contiguous_kernel(T *dst, const T *src, int64_t n) {
+    int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    int64_t stride = static_cast<int64_t>(blockDim.x) * gridDim.x;
+    for (int64_t i = idx; i < n; i += stride) {
+        dst[i] = src[i];
+    }
+}
+
+template <typename T>
+oprt_status_t launch_copy(const CopyDescriptor *desc, void *dst, const void *src, oprt_stream_t stream) {
+    constexpr int threads = 256;
+    int blocks = oprt::blocks_for(desc->elements, threads);
+    cudaStream_t s = oprt::as_cuda_stream(stream);
+    copy_contiguous_kernel<T><<<blocks, threads, 0, s>>>(
+        static_cast<T *>(dst), static_cast<const T *>(src), desc->elements);
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+} // namespace
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *dst,
+    const oprt_tensor_view_t *src) {
+    if (desc == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *desc = nullptr;
+    auto status = oprt::check_tensor(dst);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_tensor(src);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_same_dtype(*dst, *src);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_same_shape(*dst, *src);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    if (!oprt::elementwise_fast_path(*dst, *src)) {
+        return OPRT_ERR_NOT_SUPPORTED;
+    }
+
+    auto *typed = new CopyDescriptor();
+    typed->dst_view = *dst;
+    typed->src_view = *src;
+    typed->elements = oprt::numel(*dst);
+    typed->workspace_size = 0;
+    *desc = typed;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size(
+    oprt_operator_descriptor_t desc,
+    size_t *size) {
+    if (desc == nullptr || size == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *size = desc->workspace_size;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy(
+    oprt_operator_descriptor_t desc,
+    void *,
+    size_t workspace_size,
+    void *dst,
+    const void *src,
+    oprt_stream_t stream) {
+    if (desc == nullptr || dst == nullptr || src == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    if (workspace_size < desc->workspace_size) {
+        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
+    }
+    auto *typed = static_cast<const CopyDescriptor *>(desc);
+    switch (typed->dst_view.dtype) {
+    case OPRT_DTYPE_F16:
+        return launch_copy<half>(typed, dst, src, stream);
+    case OPRT_DTYPE_F32:
+        return launch_copy<float>(typed, dst, src, stream);
+    default:
+        return OPRT_ERR_UNSUPPORTED_DTYPE;
+    }
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}
--- a/ops/copy/nvidia/copy_cuda.cu
+++ b/ops/copy/nvidia/copy_cuda.cu
@ -1,9 +1,9 @@
-#include "ops/copy/nvidia/copy_cuda.h"
+#include "operator_runtime/ops/copy.h"

 #include "operator_runtime/descriptor.h"
-#include "operator_runtime/elementwise.h"
-#include "operator_runtime/tensor_checks.h"
-#include "operator_runtime/cuda_helpers.h"
+#include "operator_runtime/detail/elementwise.h"
+#include "operator_runtime/detail/tensor_checks.h"
+#include "operator_runtime/detail/cuda_helpers.h"
 #include "ops/copy/nvidia/kernel.cuh"

 #include <cuda_fp16.h>
@ -14,7 +14,6 @@ struct CopyDescriptor final : oprt_operator_descriptor {
    oprt_tensor_view_t dst_view;
    oprt_tensor_view_t src_view;
    int64_t elements = 0;
-    bool fast_path = false;

    const char *op_name() const override {
        return "copy";
@ -25,7 +24,8 @@ template <typename T>
 oprt_status_t launch_copy(const CopyDescriptor *desc, void *dst, const void *src, oprt_stream_t stream) {
    constexpr int threads = 256;
    int blocks = oprt::blocks_for(desc->elements, threads);
-    oprt::copy::nvidia::copy_contiguous_kernel<T><<<blocks, threads, 0, oprt::as_cuda_stream(stream)>>>(
+    cudaStream_t s = oprt::as_cuda_stream(stream);
+    oprt::copy::nvidia::copy_contiguous_kernel<T><<<blocks, threads, 0, s>>>(
        static_cast<T *>(dst), static_cast<const T *>(src), desc->elements);
    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
    return OPRT_SUCCESS;
@ -33,7 +33,7 @@ oprt_status_t launch_copy(const CopyDescriptor *desc, void *dst, const void *src

 } // namespace

-extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *dst,
    const oprt_tensor_view_t *src) {
@ -65,13 +65,12 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia(
    typed->dst_view = *dst;
    typed->src_view = *src;
    typed->elements = oprt::numel(*dst);
-    typed->fast_path = true;
    typed->workspace_size = 0;
    *desc = typed;
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size) {
    if (desc == nullptr || size == nullptr) {
@ -81,7 +80,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy(
    oprt_operator_descriptor_t desc,
    void *,
    size_t workspace_size,
@ -105,7 +104,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia(
    }
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor(
    oprt_operator_descriptor_t desc) {
    delete desc;
    return OPRT_SUCCESS;
--- a/ops/copy/nvidia/kernel.cuh
+++ b/ops/copy/nvidia/kernel.cuh
@ -1,16 +1,19 @@
 #pragma once

+#include <cuda_fp16.h>
 #include <stdint.h>

 namespace oprt::copy::nvidia {

 template <typename T>
 __global__ void copy_contiguous_kernel(T *dst, const T *src, int64_t n) {
-    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int64_t stride = int64_t(blockDim.x) * gridDim.x;
-    for (int64_t i = idx; i < n; i += stride) {
-        dst[i] = src[i];
-    }
+    // TODO: implement a grid-stride loop copy kernel.
+    //
+    // Suggested steps:
+    // 1. Compute the global thread index.
+    // 2. Compute the grid-wide stride.
+    // 3. Loop over i = idx; i < n; i += stride.
+    // 4. Copy src[i] to dst[i].
 }

 } // namespace oprt::copy::nvidia
--- a/ops/copy/tilelang/copy_tl.py
+++ b/ops/copy/tilelang/copy_tl.py
@ -31,8 +31,17 @@ class TileLangCopyPrepared:
    src: torch.Tensor
    kernel: object

+    def run_inputs(self, *inputs: torch.Tensor) -> torch.Tensor:
+        src = self.src if not inputs else inputs[0]
+        if len(inputs) not in (0, 1):
+            raise ValueError(f"expected 0 or 1 input tensors, got {len(inputs)}")
+        return self.kernel(src)
+
+    def run_kernel(self) -> torch.Tensor:
+        return self.run_inputs()
+
    def run(self) -> None:
-        self.out.copy_(self.kernel(self.src))
+        self.out.copy_(self.run_inputs())

    def destroy(self) -> None:
        pass
--- a/ops/copy/tilelang/kernel.py
+++ b/ops/copy/tilelang/kernel.py
@ -10,10 +10,10 @@ def copy_kernel(src, BLOCK_N: int, dtype):
    src: T.Tensor((N,), dtype)
    out = T.empty((N,), dtype)

-    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
-        T.copy(
-            src[pid_n * BLOCK_N : (pid_n + 1) * BLOCK_N],
-            out[pid_n * BLOCK_N : (pid_n + 1) * BLOCK_N],
-        )
+    # TODO: implement a tile-wise copy kernel.
+    #
+    # Suggested steps:
+    # 1. Launch one TileLang kernel over the N // BLOCK_N tiles.
+    # 2. Use T.copy to move one tile from src to out.

    return out
--- a/ops/reduce_sum/metax/README.md
+++ b/ops/reduce_sum/metax/README.md
@ -1,5 +1,4 @@
-# Reduce Sum MetaX Stub
-
-The reduce_sum ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and
-htcc/mxcc rules when MetaX hardware is available.
+# Reduce Sum MetaX Backend

+The reduce_sum MetaX backend exports the standard descriptor lifecycle with
+`_metax` symbols and is built from `ops/reduce_sum/metax/reduce_sum_metax.maca`.
--- a/ops/reduce_sum/metax/reduce_sum_metax.h
+++ b/ops/reduce_sum/metax/reduce_sum_metax.h
@ -1,6 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
-
--- a/ops/reduce_sum/metax/reduce_sum_metax.maca
+++ b/ops/reduce_sum/metax/reduce_sum_metax.maca
@ -0,0 +1,132 @@
+#include "operator_runtime/ops/reduce_sum.h"
+
+#include "operator_runtime/detail/cuda_helpers.h"
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/detail/tensor_checks.h"
+
+#include <cuda_runtime.h>
+
+namespace {
+
+struct ReduceSumDescriptor final : oprt_operator_descriptor {
+    oprt_tensor_view_t out_view;
+    oprt_tensor_view_t in_view;
+    int64_t rows = 0;
+    int64_t cols = 0;
+    int64_t axis = 1;
+
+    const char *op_name() const override {
+        return "reduce_sum";
+    }
+};
+
+__global__ void reduce_sum_rowwise_kernel(float *out, const float *in, int64_t rows, int64_t cols) {
+    int64_t row = blockIdx.x;
+    if (row >= rows) {
+        return;
+    }
+
+    const float *row_ptr = in + row * cols;
+    float partial = 0.0f;
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        partial += row_ptr[col];
+    }
+
+    extern __shared__ float shared[];
+    shared[threadIdx.x] = partial;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
+        if (threadIdx.x < stride) {
+            shared[threadIdx.x] += shared[threadIdx.x + stride];
+        }
+        __syncthreads();
+    }
+
+    if (threadIdx.x == 0) {
+        out[row] = shared[0];
+    }
+}
+
+bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in, int64_t axis) {
+    return in.dtype == OPRT_DTYPE_F32 &&
+           out.dtype == OPRT_DTYPE_F32 &&
+           in.ndim == 2 &&
+           out.ndim == 1 &&
+           axis == 1 &&
+           out.shape[0] == in.shape[0] &&
+           oprt::is_contiguous(in) &&
+           oprt::is_contiguous(out);
+}
+
+} // namespace
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *out,
+    const oprt_tensor_view_t *in,
+    int64_t axis) {
+    if (desc == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *desc = nullptr;
+    auto status = oprt::check_tensor(out);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_tensor(in);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    if (!is_rowwise_case(*out, *in, axis)) {
+        return OPRT_ERR_NOT_SUPPORTED;
+    }
+
+    auto *typed = new ReduceSumDescriptor();
+    typed->out_view = *out;
+    typed->in_view = *in;
+    typed->rows = in->shape[0];
+    typed->cols = in->shape[1];
+    typed->axis = axis;
+    typed->workspace_size = 0;
+    *desc = typed;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size(
+    oprt_operator_descriptor_t desc,
+    size_t *size) {
+    if (desc == nullptr || size == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *size = desc->workspace_size;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum(
+    oprt_operator_descriptor_t desc,
+    void *,
+    size_t workspace_size,
+    void *out,
+    const void *in,
+    oprt_stream_t stream) {
+    if (desc == nullptr || out == nullptr || in == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    if (workspace_size < desc->workspace_size) {
+        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
+    }
+    auto *typed = static_cast<const ReduceSumDescriptor *>(desc);
+    cudaStream_t s = oprt::as_cuda_stream(stream);
+    constexpr int threads = 256;
+    reduce_sum_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), s>>>(
+        static_cast<float *>(out), static_cast<const float *>(in), typed->rows, typed->cols);
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}
--- a/ops/reduce_sum/nvidia/kernel.cuh
+++ b/ops/reduce_sum/nvidia/kernel.cuh
@ -6,24 +6,14 @@
 namespace oprt::reduce_sum::nvidia {

 __global__ void reduce_sum_rowwise_kernel(float *out, const float *in, int64_t rows, int64_t cols) {
-    extern __shared__ float smem[];
-    int row = blockIdx.x;
-    float sum = 0.0f;
-    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
-        sum += in[int64_t(row) * cols + col];
-    }
-    smem[threadIdx.x] = sum;
-    __syncthreads();
-
-    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) {
-            smem[threadIdx.x] += smem[threadIdx.x + stride];
-        }
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) {
-        out[row] = smem[0];
-    }
+    // TODO: implement a row-wise reduce_sum kernel with shared memory.
+    //
+    // Suggested steps:
+    // 1. Use one block per row.
+    // 2. Let each thread accumulate a partial sum over the row.
+    // 3. Store the partial sums in shared memory.
+    // 4. Reduce shared memory with a tree reduction.
+    // 5. Let thread 0 write the final row sum to out[row].
 }

 } // namespace oprt::reduce_sum::nvidia
--- a/ops/reduce_sum/nvidia/reduce_sum_cuda.cu
+++ b/ops/reduce_sum/nvidia/reduce_sum_cuda.cu
@ -1,8 +1,8 @@
-#include "ops/reduce_sum/nvidia/reduce_sum_cuda.h"
+#include "operator_runtime/ops/reduce_sum.h"

 #include "operator_runtime/descriptor.h"
-#include "operator_runtime/tensor_checks.h"
-#include "operator_runtime/cuda_helpers.h"
+#include "operator_runtime/detail/tensor_checks.h"
+#include "operator_runtime/detail/cuda_helpers.h"
 #include "ops/reduce_sum/nvidia/kernel.cuh"

 #include <cuda_runtime.h>
@ -34,7 +34,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in

 } // namespace

-extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *out,
    const oprt_tensor_view_t *in,
@ -66,7 +66,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size) {
    if (desc == nullptr || size == nullptr) {
@ -76,7 +76,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum(
    oprt_operator_descriptor_t desc,
    void *,
    size_t workspace_size,
@ -90,14 +90,15 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia(
        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
    }
    auto *typed = static_cast<const ReduceSumDescriptor *>(desc);
+    cudaStream_t s = oprt::as_cuda_stream(stream);
    constexpr int threads = 256;
-    oprt::reduce_sum::nvidia::reduce_sum_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), oprt::as_cuda_stream(stream)>>>(
+    oprt::reduce_sum::nvidia::reduce_sum_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), s>>>(
        static_cast<float *>(out), static_cast<const float *>(in), typed->rows, typed->cols);
    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor(
    oprt_operator_descriptor_t desc) {
    delete desc;
    return OPRT_SUCCESS;
--- a/ops/reduce_sum/tilelang/kernel.py
+++ b/ops/reduce_sum/tilelang/kernel.py
@ -15,15 +15,15 @@ def reduce_sum_kernel(src, BLOCK_N: int, BLOCK_M: int):
    src: T.Tensor((N, M), dtype)
    out = T.empty((N,), dtype)

-    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
-        src_local = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
-        out_local = T.alloc_fragment((BLOCK_N,), dtype)
-        T.clear(out_local)
-
-        for m_blk in T.Serial(M // BLOCK_M):
-            T.copy(src[pid_n * BLOCK_N, m_blk * BLOCK_M], src_local, disable_tma=True)
-            T.reduce_sum(src_local, out_local, dim=1, clear=False)
-
-        T.copy(out_local, out[pid_n * BLOCK_N])
+    # TODO: implement a tiled row-wise reduce_sum kernel.
+    #
+    # Suggested steps:
+    # 1. Launch one TileLang kernel over row tiles.
+    # 2. Allocate fragments for the input tile and running output tile.
+    # 3. Clear the running output fragment before accumulation.
+    # 4. Iterate over column tiles with T.Serial(M // BLOCK_M).
+    # 5. T.copy each input tile into the fragment.
+    # 6. Call T.reduce_sum on the fragment and accumulate into the output fragment.
+    # 7. Copy the final output fragment back to global memory.

    return out
--- a/ops/reduce_sum/tilelang/reduce_sum_tl.py
+++ b/ops/reduce_sum/tilelang/reduce_sum_tl.py
@ -25,9 +25,18 @@ class TileLangReduceSumPrepared:
    dim: int = 1
    kernel: object | None = None

-    def run(self) -> None:
+    def run_inputs(self, *inputs: torch.Tensor) -> torch.Tensor:
+        if len(inputs) not in (0, 1):
+            raise ValueError(f"expected 0 or 1 input tensors, got {len(inputs)}")
+        src = self.src if not inputs else inputs[0]
        assert self.kernel is not None
-        self.out.copy_(self.kernel(self.src))
+        return self.kernel(src)
+
+    def run_kernel(self) -> torch.Tensor:
+        return self.run_inputs()
+
+    def run(self) -> None:
+        self.out.copy_(self.run_inputs())

    def destroy(self) -> None:
        pass
--- a/ops/softmax/metax/README.md
+++ b/ops/softmax/metax/README.md
@ -1,5 +1,4 @@
-# Softmax MetaX Stub
-
-The softmax ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and htcc/mxcc
-rules when MetaX hardware is available.
+# Softmax MetaX Backend

+The softmax MetaX backend exports the standard descriptor lifecycle with
+`_metax` symbols and is built from `ops/softmax/metax/softmax_metax.maca`.
--- a/ops/softmax/metax/softmax_metax.h
+++ b/ops/softmax/metax/softmax_metax.h
@ -1,6 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
-
--- a/ops/softmax/metax/softmax_metax.maca
+++ b/ops/softmax/metax/softmax_metax.maca
@ -0,0 +1,155 @@
+#include "operator_runtime/ops/softmax.h"
+
+#include "operator_runtime/detail/cuda_helpers.h"
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/detail/tensor_checks.h"
+
+#include <cuda_runtime.h>
+#include <float.h>
+#include <math.h>
+
+namespace {
+
+struct SoftmaxDescriptor final : oprt_operator_descriptor {
+    oprt_tensor_view_t out_view;
+    oprt_tensor_view_t in_view;
+    int64_t rows = 0;
+    int64_t cols = 0;
+    int64_t axis = 1;
+
+    const char *op_name() const override {
+        return "softmax";
+    }
+};
+
+__global__ void softmax_rowwise_kernel(float *out, const float *in, int64_t rows, int64_t cols) {
+    int64_t row = blockIdx.x;
+    if (row >= rows) {
+        return;
+    }
+
+    const float *row_in = in + row * cols;
+    float *row_out = out + row * cols;
+
+    float local_max = -FLT_MAX;
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        local_max = fmaxf(local_max, row_in[col]);
+    }
+
+    extern __shared__ float shared[];
+    shared[threadIdx.x] = local_max;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
+        if (threadIdx.x < stride) {
+            shared[threadIdx.x] = fmaxf(shared[threadIdx.x], shared[threadIdx.x + stride]);
+        }
+        __syncthreads();
+    }
+    float row_max = shared[0];
+
+    float local_sum = 0.0f;
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        float value = expf(row_in[col] - row_max);
+        row_out[col] = value;
+        local_sum += value;
+    }
+
+    shared[threadIdx.x] = local_sum;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
+        if (threadIdx.x < stride) {
+            shared[threadIdx.x] += shared[threadIdx.x + stride];
+        }
+        __syncthreads();
+    }
+    float row_sum = shared[0];
+
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        row_out[col] /= row_sum;
+    }
+}
+
+bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in, int64_t axis) {
+    return in.dtype == OPRT_DTYPE_F32 &&
+           out.dtype == OPRT_DTYPE_F32 &&
+           in.ndim == 2 &&
+           out.ndim == 2 &&
+           axis == 1 &&
+           oprt::same_shape(out, in) &&
+           oprt::is_contiguous(in) &&
+           oprt::is_contiguous(out);
+}
+
+} // namespace
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *out,
+    const oprt_tensor_view_t *in,
+    int64_t axis) {
+    if (desc == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *desc = nullptr;
+    auto status = oprt::check_tensor(out);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_tensor(in);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    if (!is_rowwise_case(*out, *in, axis)) {
+        return OPRT_ERR_NOT_SUPPORTED;
+    }
+
+    auto *typed = new SoftmaxDescriptor();
+    typed->out_view = *out;
+    typed->in_view = *in;
+    typed->rows = in->shape[0];
+    typed->cols = in->shape[1];
+    typed->axis = axis;
+    typed->workspace_size = 0;
+    *desc = typed;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size(
+    oprt_operator_descriptor_t desc,
+    size_t *size) {
+    if (desc == nullptr || size == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *size = desc->workspace_size;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax(
+    oprt_operator_descriptor_t desc,
+    void *,
+    size_t workspace_size,
+    void *out,
+    const void *in,
+    oprt_stream_t stream) {
+    if (desc == nullptr || out == nullptr || in == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    if (workspace_size < desc->workspace_size) {
+        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
+    }
+    auto *typed = static_cast<const SoftmaxDescriptor *>(desc);
+    cudaStream_t s = oprt::as_cuda_stream(stream);
+    constexpr int threads = 256;
+    softmax_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), s>>>(
+        static_cast<float *>(out), static_cast<const float *>(in), typed->rows, typed->cols);
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}
--- a/ops/softmax/nvidia/kernel.cuh
+++ b/ops/softmax/nvidia/kernel.cuh
@ -8,44 +8,15 @@
 namespace oprt::softmax::nvidia {

 __global__ void softmax_rowwise_kernel(float *out, const float *in, int64_t rows, int64_t cols) {
-    extern __shared__ float smem[];
-    int row = blockIdx.x;
-    float local_max = -FLT_MAX;
-    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
-        float value = in[int64_t(row) * cols + col];
-        local_max = fmaxf(local_max, value);
-    }
-    smem[threadIdx.x] = local_max;
-    __syncthreads();
-
-    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) {
-            smem[threadIdx.x] = fmaxf(smem[threadIdx.x], smem[threadIdx.x + stride]);
-        }
-        __syncthreads();
-    }
-    float row_max = smem[0];
-
-    float local_sum = 0.0f;
-    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
-        float value = expf(in[int64_t(row) * cols + col] - row_max);
-        out[int64_t(row) * cols + col] = value;
-        local_sum += value;
-    }
-    smem[threadIdx.x] = local_sum;
-    __syncthreads();
-
-    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) {
-            smem[threadIdx.x] += smem[threadIdx.x + stride];
-        }
-        __syncthreads();
-    }
-    float row_sum = smem[0];
-
-    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
-        out[int64_t(row) * cols + col] /= row_sum;
-    }
+    // TODO: implement a numerically stable row-wise softmax kernel.
+    //
+    // Suggested steps:
+    // 1. Use one block per row.
+    // 2. Reduce to find the row maximum.
+    // 3. Compute exp(x - row_max), write the temporary values to out,
+    //    and accumulate their sum.
+    // 4. Reduce to get the row sum.
+    // 5. Normalize each output element by row_sum.
 }

 } // namespace oprt::softmax::nvidia
--- a/ops/softmax/nvidia/softmax_cuda.cu
+++ b/ops/softmax/nvidia/softmax_cuda.cu
@ -1,8 +1,8 @@
-#include "ops/softmax/nvidia/softmax_cuda.h"
+#include "operator_runtime/ops/softmax.h"

 #include "operator_runtime/descriptor.h"
-#include "operator_runtime/tensor_checks.h"
-#include "operator_runtime/cuda_helpers.h"
+#include "operator_runtime/detail/tensor_checks.h"
+#include "operator_runtime/detail/cuda_helpers.h"
 #include "ops/softmax/nvidia/kernel.cuh"

 #include <cuda_runtime.h>
@ -34,7 +34,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in

 } // namespace

-extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *out,
    const oprt_tensor_view_t *in,
@ -66,7 +66,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size) {
    if (desc == nullptr || size == nullptr) {
@ -76,7 +76,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax(
    oprt_operator_descriptor_t desc,
    void *,
    size_t workspace_size,
@ -90,14 +90,15 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia(
        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
    }
    auto *typed = static_cast<const SoftmaxDescriptor *>(desc);
+    cudaStream_t s = oprt::as_cuda_stream(stream);
    constexpr int threads = 256;
-    oprt::softmax::nvidia::softmax_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), oprt::as_cuda_stream(stream)>>>(
+    oprt::softmax::nvidia::softmax_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), s>>>(
        static_cast<float *>(out), static_cast<const float *>(in), typed->rows, typed->cols);
    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor(
    oprt_operator_descriptor_t desc) {
    delete desc;
    return OPRT_SUCCESS;
--- a/ops/softmax/tilelang/kernel.py
+++ b/ops/softmax/tilelang/kernel.py
@ -16,34 +16,21 @@ def softmax_kernel(src, BLOCK_N: int, BLOCK_M: int):
    src: T.Tensor((N, M), dtype)
    out = T.empty((N, M), dtype)

-    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
-        src_local = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
-        out_local = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
-        cur_exp = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
-        cur_max = T.alloc_fragment((BLOCK_N,), dtype)
-        cur_sum = T.alloc_fragment((BLOCK_N,), dtype)
-        lse = T.alloc_fragment((BLOCK_N,), dtype)
-
-        T.fill(lse, -T.infinity(dtype))
-
-        for m_blk in T.Serial(M // BLOCK_M):
-            T.copy(src[pid_n * BLOCK_N, m_blk * BLOCK_M], src_local, disable_tma=True)
-            T.reduce_max(src_local, cur_max, dim=1, clear=True)
-
-            for i, j in T.Parallel(BLOCK_N, BLOCK_M):
-                cur_exp[i, j] = T.exp2(src_local[i, j] * log2_e - cur_max[i] * log2_e)
-
-            T.reduce_sum(cur_exp, cur_sum, dim=1, clear=True)
-
-            for i in T.Parallel(BLOCK_N):
-                lse[i] = cur_max[i] * log2_e + T.log2(
-                    T.exp2(lse[i] - cur_max[i] * log2_e) + cur_sum[i]
-                )
-
-        for m_blk in T.Serial(M // BLOCK_M):
-            T.copy(src[pid_n * BLOCK_N, m_blk * BLOCK_M], src_local, disable_tma=True)
-            for i, j in T.Parallel(BLOCK_N, BLOCK_M):
-                out_local[i, j] = T.exp2(src_local[i, j] * log2_e - lse[i])
-            T.copy(out_local, out[pid_n * BLOCK_N, m_blk * BLOCK_M])
+    # TODO: implement a tiled row-wise softmax kernel.
+    #
+    # Suggested steps:
+    # 1. Launch one TileLang kernel over row tiles.
+    # 2. Allocate fragments for src, out, temporary exp values, row max, row sum, and lse.
+    # 3. Initialize the running log-sum-exp state.
+    # 4. In a first T.Serial loop over column tiles:
+    #    - copy the input tile into a fragment
+    #    - reduce to get the tile max
+    #    - compute exp2-based temporary values
+    #    - reduce to get the tile sum
+    #    - update the running lse
+    # 5. In a second T.Serial loop over column tiles:
+    #    - copy the input tile again
+    #    - normalize with the final lse
+    #    - copy the result tile to global memory

    return out
--- a/ops/softmax/tilelang/softmax_tl.py
+++ b/ops/softmax/tilelang/softmax_tl.py
@ -25,9 +25,18 @@ class TileLangSoftmaxPrepared:
    dim: int = 1
    kernel: object | None = None

-    def run(self) -> None:
+    def run_inputs(self, *inputs: torch.Tensor) -> torch.Tensor:
+        if len(inputs) not in (0, 1):
+            raise ValueError(f"expected 0 or 1 input tensors, got {len(inputs)}")
+        src = self.src if not inputs else inputs[0]
        assert self.kernel is not None
-        self.out.copy_(self.kernel(self.src))
+        return self.kernel(src)
+
+    def run_kernel(self) -> torch.Tensor:
+        return self.run_inputs()
+
+    def run(self) -> None:
+        self.out.copy_(self.run_inputs())

    def destroy(self) -> None:
        pass
--- a/ops/vector_add/metax/README.md
+++ b/ops/vector_add/metax/README.md
@ -1,5 +1,4 @@
-# Vector Add MetaX Stub
-
-The vector_add ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and
-htcc/mxcc rules when MetaX hardware is available.
+# Vector Add MetaX Backend

+The vector_add MetaX backend exports the standard descriptor lifecycle with
+`_metax` symbols and is built from `ops/vector_add/metax/vector_add_metax.maca`.
--- a/ops/vector_add/metax/vector_add_metax.h
+++ b/ops/vector_add/metax/vector_add_metax.h
@ -1,6 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
-
--- a/ops/vector_add/metax/vector_add_metax.maca
+++ b/ops/vector_add/metax/vector_add_metax.maca
@ -0,0 +1,141 @@
+#include "operator_runtime/ops/vector_add.h"
+
+#include "operator_runtime/detail/cuda_helpers.h"
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/detail/elementwise.h"
+#include "operator_runtime/detail/tensor_checks.h"
+
+#include <cuda_fp16.h>
+
+namespace {
+
+struct VectorAddDescriptor final : oprt_operator_descriptor {
+    oprt_tensor_view_t out_view;
+    oprt_tensor_view_t a_view;
+    oprt_tensor_view_t b_view;
+    int64_t elements = 0;
+
+    const char *op_name() const override {
+        return "vector_add";
+    }
+};
+
+template <typename T>
+__device__ T add_values(T a, T b) {
+    return a + b;
+}
+
+template <>
+__device__ inline half add_values<half>(half a, half b) {
+    return __hadd(a, b);
+}
+
+template <typename T>
+__global__ void vector_add_contiguous_kernel(
+    T *__restrict__ out,
+    const T *__restrict__ a,
+    const T *__restrict__ b,
+    int64_t n) {
+    int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    int64_t stride = static_cast<int64_t>(blockDim.x) * gridDim.x;
+    for (int64_t i = idx; i < n; i += stride) {
+        out[i] = add_values(a[i], b[i]);
+    }
+}
+
+template <typename T>
+oprt_status_t launch_vector_add(const VectorAddDescriptor *desc, void *out, const void *a, const void *b, oprt_stream_t stream) {
+    constexpr int threads = 256;
+    int blocks = oprt::blocks_for(desc->elements, threads);
+    cudaStream_t s = oprt::as_cuda_stream(stream);
+    vector_add_contiguous_kernel<T><<<blocks, threads, 0, s>>>(
+        static_cast<T *>(out), static_cast<const T *>(a), static_cast<const T *>(b), desc->elements);
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+} // namespace
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *out,
+    const oprt_tensor_view_t *a,
+    const oprt_tensor_view_t *b) {
+    if (desc == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *desc = nullptr;
+    auto status = oprt::check_tensor(out);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_tensor(a);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_tensor(b);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_same_dtype(*out, *a);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    status = oprt::check_same_dtype(*out, *b);
+    if (status != OPRT_SUCCESS) {
+        return status;
+    }
+    if (!oprt::elementwise_fast_path(*out, *a, *b)) {
+        return OPRT_ERR_NOT_SUPPORTED;
+    }
+
+    auto *typed = new VectorAddDescriptor();
+    typed->out_view = *out;
+    typed->a_view = *a;
+    typed->b_view = *b;
+    typed->elements = oprt::numel(*out);
+    typed->workspace_size = 0;
+    *desc = typed;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size(
+    oprt_operator_descriptor_t desc,
+    size_t *size) {
+    if (desc == nullptr || size == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    *size = desc->workspace_size;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add(
+    oprt_operator_descriptor_t desc,
+    void *,
+    size_t workspace_size,
+    void *out,
+    const void *a,
+    const void *b,
+    oprt_stream_t stream) {
+    if (desc == nullptr || out == nullptr || a == nullptr || b == nullptr) {
+        return OPRT_ERR_INVALID_ARG;
+    }
+    if (workspace_size < desc->workspace_size) {
+        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
+    }
+    auto *typed = static_cast<const VectorAddDescriptor *>(desc);
+    switch (typed->out_view.dtype) {
+    case OPRT_DTYPE_F16:
+        return launch_vector_add<half>(typed, out, a, b, stream);
+    case OPRT_DTYPE_F32:
+        return launch_vector_add<float>(typed, out, a, b, stream);
+    default:
+        return OPRT_ERR_UNSUPPORTED_DTYPE;
+    }
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}
--- a/ops/vector_add/nvidia/kernel.cuh
+++ b/ops/vector_add/nvidia/kernel.cuh
@ -7,21 +7,28 @@ namespace oprt::vector_add::nvidia {

 template <typename T>
 __device__ T add_values(T a, T b) {
-    return a + b;
+    // TODO: return the elementwise sum for generic types.
+    return T{};
 }

 template <>
 __device__ inline half add_values<half>(half a, half b) {
-    return __hadd(a, b);
+    // TODO: return the half-precision elementwise sum.
+    return half{};
 }

 template <typename T>
-__global__ void vector_add_contiguous_kernel(T *out, const T *a, const T *b, int64_t n) {
-    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
-    int64_t stride = int64_t(blockDim.x) * gridDim.x;
-    for (int64_t i = idx; i < n; i += stride) {
-        out[i] = add_values(a[i], b[i]);
-    }
+__global__ void vector_add_contiguous_kernel(T * __restrict__ out,
+                                             const T * __restrict__ a,
+                                             const T * __restrict__ b,
+                                             int64_t n) {
+    // TODO: implement a grid-stride loop vector add kernel.
+    //
+    // Suggested steps:
+    // 1. Compute the global thread index.
+    // 2. Compute the grid-wide stride.
+    // 3. Loop over i = idx; i < n; i += stride.
+    // 4. Write out[i] = add_values(a[i], b[i]).
 }

 } // namespace oprt::vector_add::nvidia
--- a/ops/vector_add/nvidia/vector_add_cuda.cu
+++ b/ops/vector_add/nvidia/vector_add_cuda.cu
@ -1,9 +1,9 @@
-#include "ops/vector_add/nvidia/vector_add_cuda.h"
+#include "operator_runtime/ops/vector_add.h"

 #include "operator_runtime/descriptor.h"
-#include "operator_runtime/elementwise.h"
-#include "operator_runtime/tensor_checks.h"
-#include "operator_runtime/cuda_helpers.h"
+#include "operator_runtime/detail/elementwise.h"
+#include "operator_runtime/detail/tensor_checks.h"
+#include "operator_runtime/detail/cuda_helpers.h"
 #include "ops/vector_add/nvidia/kernel.cuh"

 #include <cuda_fp16.h>
@ -25,7 +25,8 @@ template <typename T>
 oprt_status_t launch_vector_add(const VectorAddDescriptor *desc, void *out, const void *a, const void *b, oprt_stream_t stream) {
    constexpr int threads = 256;
    int blocks = oprt::blocks_for(desc->elements, threads);
-    oprt::vector_add::nvidia::vector_add_contiguous_kernel<T><<<blocks, threads, 0, oprt::as_cuda_stream(stream)>>>(
+    cudaStream_t s = oprt::as_cuda_stream(stream);
+    oprt::vector_add::nvidia::vector_add_contiguous_kernel<T><<<blocks, threads, 0, s>>>(
        static_cast<T *>(out), static_cast<const T *>(a), static_cast<const T *>(b), desc->elements);
    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
    return OPRT_SUCCESS;
@ -33,7 +34,7 @@ oprt_status_t launch_vector_add(const VectorAddDescriptor *desc, void *out, cons

 } // namespace

-extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor(
    oprt_operator_descriptor_t *desc,
    const oprt_tensor_view_t *out,
    const oprt_tensor_view_t *a,
@ -76,7 +77,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size(
    oprt_operator_descriptor_t desc,
    size_t *size) {
    if (desc == nullptr || size == nullptr) {
@ -86,7 +87,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_nvidia(
    return OPRT_SUCCESS;
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add(
    oprt_operator_descriptor_t desc,
    void *,
    size_t workspace_size,
@ -111,7 +112,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia(
    }
 }

-extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor_nvidia(
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor(
    oprt_operator_descriptor_t desc) {
    delete desc;
    return OPRT_SUCCESS;
--- a/ops/vector_add/tilelang/kernel.py
+++ b/ops/vector_add/tilelang/kernel.py
@ -11,9 +11,11 @@ def vector_add_kernel(a, b, BLOCK_N: int, dtype):
    b: T.Tensor((N,), dtype)
    out = T.empty((N,), dtype)

-    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
-        base = pid_n * BLOCK_N
-        for i in T.Parallel(BLOCK_N):
-            out[base + i] = a[base + i] + b[base + i]
+    # TODO: implement a tile-wise vector add kernel.
+    #
+    # Suggested steps:
+    # 1. Launch one TileLang kernel over the N // BLOCK_N tiles.
+    # 2. Compute the tile base offset.
+    # 3. Use T.Parallel(BLOCK_N) to fill out[base + i] = a[base + i] + b[base + i].

    return out
--- a/ops/vector_add/tilelang/vector_add_tl.py
+++ b/ops/vector_add/tilelang/vector_add_tl.py
@ -32,8 +32,17 @@ class TileLangVectorAddPrepared:
    b: torch.Tensor
    kernel: object

+    def run_inputs(self, *inputs: torch.Tensor) -> torch.Tensor:
+        if len(inputs) not in (0, 2):
+            raise ValueError(f"expected 0 or 2 input tensors, got {len(inputs)}")
+        a, b = (self.a, self.b) if not inputs else inputs
+        return self.kernel(a, b)
+
+    def run_kernel(self) -> torch.Tensor:
+        return self.run_inputs()
+
    def run(self) -> None:
-        self.out.copy_(self.kernel(self.a, self.b))
+        self.out.copy_(self.run_inputs())

    def destroy(self) -> None:
        pass
--- a/python/operator_runtime/_internal/init.py
+++ b/python/operator_runtime/_internal/init.py
@ -1,6 +1,6 @@
 from .loader import load_library
 from .tensor_view import TensorView, tensor_view, dtype_to_oprt, current_stream_ptr, OPRT_MAX_DIMS
-from .ctypes_bindings import (
+from .bindings import (
    CFunctions,
    Descriptor,
    OperatorRuntimeError,
--- a/python/operator_runtime/_internal/bindings.py
+++ b/python/operator_runtime/_internal/bindings.py
@ -0,0 +1,181 @@
+from __future__ import annotations
+
+import ctypes
+from dataclasses import dataclass
+from typing import Callable
+
+from operator_runtime.backend import Backend
+
+from .loader import load_library
+from .tensor_view import TensorView
+
+Status = ctypes.c_int
+Descriptor = ctypes.c_void_p
+
+
+class OperatorRuntimeError(RuntimeError):
+    pass
+
+
+def check_status(status: int) -> None:
+    if status == 0:
+        return
+    lib = load_library()
+    lib.oprt_status_string.argtypes = [ctypes.c_int]
+    lib.oprt_status_string.restype = ctypes.c_char_p
+    message = lib.oprt_status_string(status).decode("utf-8")
+    raise OperatorRuntimeError(message)
+
+
+@dataclass(frozen=True)
+class CFunctions:
+    create: Callable
+    workspace: Callable
+    execute: Callable
+    destroy: Callable
+
+
+def _missing_symbol_error(name: str, symbol: str, exc: AttributeError) -> OperatorRuntimeError:
+    raise OperatorRuntimeError(
+        f"missing symbol {symbol} for operator {name}"
+    ) from exc
+
+
+def bind_unary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions:
+    lib = load_library()
+    create_symbol = f"oprt_create_{name}_descriptor"
+    try:
+        create = getattr(lib, create_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, create_symbol, exc)
+    create.argtypes = [ctypes.POINTER(Descriptor), ctypes.POINTER(TensorView), ctypes.POINTER(TensorView)]
+    create.restype = Status
+
+    workspace_symbol = f"oprt_get_{name}_workspace_size"
+    try:
+        workspace = getattr(lib, workspace_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, workspace_symbol, exc)
+    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
+    workspace.restype = Status
+
+    execute_symbol = f"oprt_execute_{name}"
+    try:
+        execute = getattr(lib, execute_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, execute_symbol, exc)
+    execute.argtypes = [
+        Descriptor,
+        ctypes.c_void_p,
+        ctypes.c_size_t,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+    ]
+    execute.restype = Status
+
+    destroy_symbol = f"oprt_destroy_{name}_descriptor"
+    try:
+        destroy = getattr(lib, destroy_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, destroy_symbol, exc)
+    destroy.argtypes = [Descriptor]
+    destroy.restype = Status
+    return CFunctions(create, workspace, execute, destroy)
+
+
+def bind_binary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions:
+    lib = load_library()
+    create_symbol = f"oprt_create_{name}_descriptor"
+    try:
+        create = getattr(lib, create_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, create_symbol, exc)
+    create.argtypes = [
+        ctypes.POINTER(Descriptor),
+        ctypes.POINTER(TensorView),
+        ctypes.POINTER(TensorView),
+        ctypes.POINTER(TensorView),
+    ]
+    create.restype = Status
+
+    workspace_symbol = f"oprt_get_{name}_workspace_size"
+    try:
+        workspace = getattr(lib, workspace_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, workspace_symbol, exc)
+    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
+    workspace.restype = Status
+
+    execute_symbol = f"oprt_execute_{name}"
+    try:
+        execute = getattr(lib, execute_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, execute_symbol, exc)
+    execute.argtypes = [
+        Descriptor,
+        ctypes.c_void_p,
+        ctypes.c_size_t,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+    ]
+    execute.restype = Status
+
+    destroy_symbol = f"oprt_destroy_{name}_descriptor"
+    try:
+        destroy = getattr(lib, destroy_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, destroy_symbol, exc)
+    destroy.argtypes = [Descriptor]
+    destroy.restype = Status
+    return CFunctions(create, workspace, execute, destroy)
+
+
+def bind_reduce_like(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions:
+    lib = load_library()
+    create_symbol = f"oprt_create_{name}_descriptor"
+    try:
+        create = getattr(lib, create_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, create_symbol, exc)
+    create.argtypes = [
+        ctypes.POINTER(Descriptor),
+        ctypes.POINTER(TensorView),
+        ctypes.POINTER(TensorView),
+        ctypes.c_int64,
+    ]
+    create.restype = Status
+
+    workspace_symbol = f"oprt_get_{name}_workspace_size"
+    try:
+        workspace = getattr(lib, workspace_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, workspace_symbol, exc)
+    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
+    workspace.restype = Status
+
+    execute_symbol = f"oprt_execute_{name}"
+    try:
+        execute = getattr(lib, execute_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, execute_symbol, exc)
+    execute.argtypes = [
+        Descriptor,
+        ctypes.c_void_p,
+        ctypes.c_size_t,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+    ]
+    execute.restype = Status
+
+    destroy_symbol = f"oprt_destroy_{name}_descriptor"
+    try:
+        destroy = getattr(lib, destroy_symbol)
+    except AttributeError as exc:
+        _missing_symbol_error(name, destroy_symbol, exc)
+    destroy.argtypes = [Descriptor]
+    destroy.restype = Status
+    return CFunctions(create, workspace, execute, destroy)
--- a/python/operator_runtime/_internal/loader.py
+++ b/python/operator_runtime/_internal/loader.py
@ -12,11 +12,13 @@ def _candidate_library_paths() -> list[Path]:
    if build_dir:
        root = Path(build_dir)
        candidates.extend([root / "libcamp_ops.so", root / "ops" / "libcamp_ops.so"])
-    repo_root = Path(__file__).resolve().parents[2]
+    repo_root = Path(__file__).resolve().parents[3]
    candidates.extend(
        [
            repo_root / "build" / "libcamp_ops.so",
            repo_root / "build" / "ops" / "libcamp_ops.so",
+            repo_root / "build-metax" / "libcamp_ops.so",
+            repo_root / "build-metax" / "ops" / "libcamp_ops.so",
        ]
    )
    return candidates
@ -29,4 +31,3 @@ def load_library() -> ctypes.CDLL:
            return ctypes.CDLL(str(path))
    searched = ", ".join(str(path) for path in _candidate_library_paths())
    raise FileNotFoundError(f"libcamp_ops.so not found; searched: {searched}")
-
--- a/python/operator_runtime/_internal/prepared.py
+++ b/python/operator_runtime/_internal/prepared.py
@ -6,7 +6,7 @@ from typing import Any

 import torch

-from .ctypes_bindings import CFunctions, Descriptor, check_status
+from .bindings import CFunctions, Descriptor, check_status


@dataclass
@ -18,6 +18,20 @@ class PreparedOp:
    stream_tensor: torch.Tensor | None = None

    def run(self) -> None:
+        self.run_inputs()
+
+    def run_inputs(self, *runner_tensors: torch.Tensor) -> None:
+        if runner_tensors:
+            expected = len(self.runner_args) - 1
+            if len(runner_tensors) != expected:
+                raise ValueError(f"expected {expected} input tensors, got {len(runner_tensors)}")
+            runner_args = (
+                self.runner_args[0],
+                *tuple(ctypes.c_void_p(tensor.data_ptr()) for tensor in runner_tensors),
+            )
+        else:
+            runner_args = self.runner_args
+
        stream = torch.cuda.current_stream(device=self.stream_tensor.device if self.stream_tensor is not None else None)
        workspace_ptr = None if self.workspace is None else ctypes.c_void_p(self.workspace.data_ptr())
        workspace_size = 0 if self.workspace is None else self.workspace.numel()
@ -25,11 +39,14 @@ class PreparedOp:
            self.descriptor,
            workspace_ptr,
            workspace_size,
-            *self.runner_args,
+            *runner_args,
            ctypes.c_void_p(stream.cuda_stream),
        )
        check_status(status)

+    def run_kernel(self) -> None:
+        self.run()
+
    def destroy(self) -> None:
        if self.descriptor:
            check_status(self.funcs.destroy(self.descriptor))
@ -46,4 +63,3 @@ class PreparedOp:
            self.destroy()
        except Exception:
            pass
-
--- a/python/operator_runtime/_internal/tensor_view.py
+++ b/python/operator_runtime/_internal/tensor_view.py
@ -52,4 +52,3 @@ def current_stream_ptr(tensor: torch.Tensor | None = None) -> ctypes.c_void_p:
    device = tensor.device if tensor is not None and tensor.is_cuda else None
    stream = torch.cuda.current_stream(device=device)
    return ctypes.c_void_p(stream.cuda_stream)
-
--- a/python/operator_runtime/_runtime/ctypes_bindings.py
+++ b/python/operator_runtime/_runtime/ctypes_bindings.py
@ -1,126 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-from dataclasses import dataclass
-from typing import Callable
-
-from .loader import load_library
-from .tensor_view import TensorView
-
-Status = ctypes.c_int
-Descriptor = ctypes.c_void_p
-
-
-class OperatorRuntimeError(RuntimeError):
-    pass
-
-
-def check_status(status: int) -> None:
-    if status == 0:
-        return
-    lib = load_library()
-    lib.oprt_status_string.argtypes = [ctypes.c_int]
-    lib.oprt_status_string.restype = ctypes.c_char_p
-    message = lib.oprt_status_string(status).decode("utf-8")
-    raise OperatorRuntimeError(message)
-
-
-@dataclass(frozen=True)
-class CFunctions:
-    create: Callable
-    workspace: Callable
-    execute: Callable
-    destroy: Callable
-
-
-def bind_unary(name: str) -> CFunctions:
-    lib = load_library()
-    create = getattr(lib, f"oprt_create_{name}_descriptor_nvidia")
-    create.argtypes = [ctypes.POINTER(Descriptor), ctypes.POINTER(TensorView), ctypes.POINTER(TensorView)]
-    create.restype = Status
-
-    workspace = getattr(lib, f"oprt_get_{name}_workspace_size_nvidia")
-    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
-    workspace.restype = Status
-
-    execute = getattr(lib, f"oprt_execute_{name}_nvidia")
-    execute.argtypes = [
-        Descriptor,
-        ctypes.c_void_p,
-        ctypes.c_size_t,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-    ]
-    execute.restype = Status
-
-    destroy = getattr(lib, f"oprt_destroy_{name}_descriptor_nvidia")
-    destroy.argtypes = [Descriptor]
-    destroy.restype = Status
-    return CFunctions(create, workspace, execute, destroy)
-
-
-def bind_binary(name: str) -> CFunctions:
-    lib = load_library()
-    create = getattr(lib, f"oprt_create_{name}_descriptor_nvidia")
-    create.argtypes = [
-        ctypes.POINTER(Descriptor),
-        ctypes.POINTER(TensorView),
-        ctypes.POINTER(TensorView),
-        ctypes.POINTER(TensorView),
-    ]
-    create.restype = Status
-
-    workspace = getattr(lib, f"oprt_get_{name}_workspace_size_nvidia")
-    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
-    workspace.restype = Status
-
-    execute = getattr(lib, f"oprt_execute_{name}_nvidia")
-    execute.argtypes = [
-        Descriptor,
-        ctypes.c_void_p,
-        ctypes.c_size_t,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-    ]
-    execute.restype = Status
-
-    destroy = getattr(lib, f"oprt_destroy_{name}_descriptor_nvidia")
-    destroy.argtypes = [Descriptor]
-    destroy.restype = Status
-    return CFunctions(create, workspace, execute, destroy)
-
-
-def bind_reduce_like(name: str) -> CFunctions:
-    lib = load_library()
-    create = getattr(lib, f"oprt_create_{name}_descriptor_nvidia")
-    create.argtypes = [
-        ctypes.POINTER(Descriptor),
-        ctypes.POINTER(TensorView),
-        ctypes.POINTER(TensorView),
-        ctypes.c_int64,
-    ]
-    create.restype = Status
-
-    workspace = getattr(lib, f"oprt_get_{name}_workspace_size_nvidia")
-    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
-    workspace.restype = Status
-
-    execute = getattr(lib, f"oprt_execute_{name}_nvidia")
-    execute.argtypes = [
-        Descriptor,
-        ctypes.c_void_p,
-        ctypes.c_size_t,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-        ctypes.c_void_p,
-    ]
-    execute.restype = Status
-
-    destroy = getattr(lib, f"oprt_destroy_{name}_descriptor_nvidia")
-    destroy.argtypes = [Descriptor]
-    destroy.restype = Status
-    return CFunctions(create, workspace, execute, destroy)
-
--- a/python/operator_runtime/ops/_common.py
+++ b/python/operator_runtime/ops/_common.py
@ -0,0 +1,24 @@
+from __future__ import annotations
+
+import ctypes
+
+import torch
+
+from operator_runtime._internal import CFunctions, Descriptor, PreparedOp, check_status
+
+
+def build_prepared_op(
+    funcs: CFunctions,
+    create_args: tuple[object, ...],
+    tensors: tuple[torch.Tensor, ...],
+    stream_tensor: torch.Tensor,
+) -> PreparedOp:
+    desc = Descriptor()
+    check_status(funcs.create(ctypes.byref(desc), *create_args))
+    workspace_size = ctypes.c_size_t()
+    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
+    workspace = None
+    if workspace_size.value:
+        workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=stream_tensor.device)
+    runner_args = tuple(ctypes.c_void_p(tensor.data_ptr()) for tensor in tensors)
+    return PreparedOp(funcs, desc, workspace, runner_args, stream_tensor)
--- a/python/operator_runtime/ops/copy.py
+++ b/python/operator_runtime/ops/copy.py
@ -5,7 +5,8 @@ import ctypes
 import torch

 from operator_runtime.backend import Backend, normalize_backend
-from operator_runtime._runtime import Descriptor, bind_unary, check_status, PreparedOp, tensor_view
+from operator_runtime._internal import PreparedOp, bind_unary, tensor_view
+from operator_runtime.ops._common import build_prepared_op


 def _check(out: torch.Tensor, src: torch.Tensor) -> None:
@ -26,19 +27,14 @@ def prepare_copy(out: torch.Tensor, src: torch.Tensor, backend: str | Backend =
        from ops.copy.tilelang.copy_tl import prepare_copy_tl

        return prepare_copy_tl(out, src)
-    if backend is not Backend.NVIDIA:
+    if backend not in (Backend.NVIDIA, Backend.METAX):
        raise NotImplementedError(f"backend {backend.value} is not runnable")

-    funcs = bind_unary("copy")
-    desc = Descriptor()
+    funcs = bind_unary("copy", backend)
    out_view = tensor_view(out)
    src_view = tensor_view(src)
-    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(src_view)))
-    workspace_size = ctypes.c_size_t()
-    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
-    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
-    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(src.data_ptr()))
-    return PreparedOp(funcs, desc, workspace, args, out)
+    create_args = (ctypes.byref(out_view), ctypes.byref(src_view))
+    return build_prepared_op(funcs, create_args, (out, src), out)


 def copy_(out: torch.Tensor, src: torch.Tensor, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
@ -50,4 +46,3 @@ def copy_(out: torch.Tensor, src: torch.Tensor, backend: str | Backend = Backend
 def copy(src: torch.Tensor, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
    out = torch.empty_like(src)
    return copy_(out, src, backend)
-
--- a/python/operator_runtime/ops/reduce_sum.py
+++ b/python/operator_runtime/ops/reduce_sum.py
@ -5,7 +5,8 @@ import ctypes
 import torch

 from operator_runtime.backend import Backend, normalize_backend
-from operator_runtime._runtime import Descriptor, bind_reduce_like, check_status, PreparedOp, tensor_view
+from operator_runtime._internal import PreparedOp, bind_reduce_like, tensor_view
+from operator_runtime.ops._common import build_prepared_op


 def _check(out: torch.Tensor, src: torch.Tensor, dim: int) -> None:
@ -33,19 +34,14 @@ def prepare_reduce_sum(
        from ops.reduce_sum.tilelang.reduce_sum_tl import prepare_reduce_sum_tl

        return prepare_reduce_sum_tl(out, src, dim=dim)
-    if backend is not Backend.NVIDIA:
+    if backend not in (Backend.NVIDIA, Backend.METAX):
        raise NotImplementedError(f"backend {backend.value} is not runnable")

-    funcs = bind_reduce_like("reduce_sum")
-    desc = Descriptor()
+    funcs = bind_reduce_like("reduce_sum", backend)
    out_view = tensor_view(out)
    src_view = tensor_view(src)
-    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(src_view), ctypes.c_int64(dim)))
-    workspace_size = ctypes.c_size_t()
-    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
-    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
-    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(src.data_ptr()))
-    return PreparedOp(funcs, desc, workspace, args, out)
+    create_args = (ctypes.byref(out_view), ctypes.byref(src_view), ctypes.c_int64(dim))
+    return build_prepared_op(funcs, create_args, (out, src), out)


 def reduce_sum_(out: torch.Tensor, src: torch.Tensor, dim: int = 1, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
@ -59,4 +55,3 @@ def reduce_sum(src: torch.Tensor, dim: int = 1, backend: str | Backend = Backend
        raise ValueError("reduce_sum v1 supports 2D row-wise reduction over dim=1")
    out = torch.empty((src.shape[0],), dtype=src.dtype, device=src.device)
    return reduce_sum_(out, src, dim, backend)
-
--- a/python/operator_runtime/ops/softmax.py
+++ b/python/operator_runtime/ops/softmax.py
@ -5,7 +5,8 @@ import ctypes
 import torch

 from operator_runtime.backend import Backend, normalize_backend
-from operator_runtime._runtime import Descriptor, bind_reduce_like, check_status, PreparedOp, tensor_view
+from operator_runtime._internal import PreparedOp, bind_reduce_like, tensor_view
+from operator_runtime.ops._common import build_prepared_op


 def _check(out: torch.Tensor, src: torch.Tensor, dim: int) -> None:
@ -33,19 +34,14 @@ def prepare_softmax(
        from ops.softmax.tilelang.softmax_tl import prepare_softmax_tl

        return prepare_softmax_tl(out, src, dim=dim)
-    if backend is not Backend.NVIDIA:
+    if backend not in (Backend.NVIDIA, Backend.METAX):
        raise NotImplementedError(f"backend {backend.value} is not runnable")

-    funcs = bind_reduce_like("softmax")
-    desc = Descriptor()
+    funcs = bind_reduce_like("softmax", backend)
    out_view = tensor_view(out)
    src_view = tensor_view(src)
-    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(src_view), ctypes.c_int64(dim)))
-    workspace_size = ctypes.c_size_t()
-    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
-    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
-    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(src.data_ptr()))
-    return PreparedOp(funcs, desc, workspace, args, out)
+    create_args = (ctypes.byref(out_view), ctypes.byref(src_view), ctypes.c_int64(dim))
+    return build_prepared_op(funcs, create_args, (out, src), out)


 def softmax_(out: torch.Tensor, src: torch.Tensor, dim: int = 1, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
@ -57,4 +53,3 @@ def softmax_(out: torch.Tensor, src: torch.Tensor, dim: int = 1, backend: str |
 def softmax(src: torch.Tensor, dim: int = 1, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
    out = torch.empty_like(src)
    return softmax_(out, src, dim, backend)
-
--- a/python/operator_runtime/ops/vector_add.py
+++ b/python/operator_runtime/ops/vector_add.py
@ -5,7 +5,8 @@ import ctypes
 import torch

 from operator_runtime.backend import Backend, normalize_backend
-from operator_runtime._runtime import Descriptor, bind_binary, check_status, PreparedOp, tensor_view
+from operator_runtime._internal import PreparedOp, bind_binary, tensor_view
+from operator_runtime.ops._common import build_prepared_op


 def _check(out: torch.Tensor, a: torch.Tensor, b: torch.Tensor) -> None:
@ -31,20 +32,15 @@ def prepare_vector_add(
        from ops.vector_add.tilelang.vector_add_tl import prepare_vector_add_tl

        return prepare_vector_add_tl(out, a, b)
-    if backend is not Backend.NVIDIA:
+    if backend not in (Backend.NVIDIA, Backend.METAX):
        raise NotImplementedError(f"backend {backend.value} is not runnable")

-    funcs = bind_binary("vector_add")
-    desc = Descriptor()
+    funcs = bind_binary("vector_add", backend)
    out_view = tensor_view(out)
    a_view = tensor_view(a)
    b_view = tensor_view(b)
-    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(a_view), ctypes.byref(b_view)))
-    workspace_size = ctypes.c_size_t()
-    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
-    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
-    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(a.data_ptr()), ctypes.c_void_p(b.data_ptr()))
-    return PreparedOp(funcs, desc, workspace, args, out)
+    create_args = (ctypes.byref(out_view), ctypes.byref(a_view), ctypes.byref(b_view))
+    return build_prepared_op(funcs, create_args, (out, a, b), out)


 def vector_add_(out: torch.Tensor, a: torch.Tensor, b: torch.Tensor, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
@ -56,4 +52,3 @@ def vector_add_(out: torch.Tensor, a: torch.Tensor, b: torch.Tensor, backend: st
 def vector_add(a: torch.Tensor, b: torch.Tensor, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
    out = torch.empty_like(a)
    return vector_add_(out, a, b, backend)
-
--- a/python/operator_runtime/testing/benchmark.py
+++ b/python/operator_runtime/testing/benchmark.py
@ -1,19 +0,0 @@
-from __future__ import annotations
-
-from collections.abc import Callable
-
-import torch
-
-
-def cuda_time_ms(fn: Callable[[], None], *, warmup: int = 10, iterations: int = 100) -> float:
-    for _ in range(warmup):
-        fn()
-    torch.cuda.synchronize()
-    start = torch.cuda.Event(enable_timing=True)
-    end = torch.cuda.Event(enable_timing=True)
-    start.record()
-    for _ in range(iterations):
-        fn()
-    end.record()
-    torch.cuda.synchronize()
-    return start.elapsed_time(end) / iterations
--- a/python/operator_runtime_testing/init.py
+++ b/python/operator_runtime_testing/init.py
--- a/python/operator_runtime_testing/assertions.py
+++ b/python/operator_runtime_testing/assertions.py
--- a/python/operator_runtime_testing/benchmark.py
+++ b/python/operator_runtime_testing/benchmark.py
@ -0,0 +1,136 @@
+from __future__ import annotations
+
+from collections.abc import Callable
+
+import torch
+from torch.autograd.profiler import DeviceType
+
+_l2_flush_cache: torch.Tensor | None = None
+
+
+def _get_l2_flush_cache() -> torch.Tensor:
+    global _l2_flush_cache
+    if _l2_flush_cache is None:
+        l2_bytes = torch.cuda.get_device_properties(0).L2_cache_size
+        if l2_bytes <= 0:
+            l2_bytes = int(256e6)
+        _l2_flush_cache = torch.empty(max(l2_bytes // 4, 1), dtype=torch.int, device="cuda")
+    return _l2_flush_cache
+
+
+def _sum_cuda_kernel_ms(kineto_results) -> float:
+    total_us = 0.0
+    for evt in kineto_results.events():
+        if evt.device_type() != DeviceType.CUDA:
+            continue
+        name = evt.name()
+        if "vectorized_elementwise" in name and "FillFunctor" in name:
+            continue
+        total_us += evt.duration_ns() / 1000.0
+    return total_us / 1000.0
+
+
+def _build_arg_pool(
+    args: tuple[object, ...],
+    *,
+    pool_size: int,
+    max_clone_bytes: int,
+) -> list[tuple[object, ...]] | None:
+    if not args:
+        return None
+
+    tensor_mask = tuple(isinstance(arg, torch.Tensor) for arg in args)
+    total_bytes = sum(
+        arg.nelement() * arg.element_size()
+        for arg, is_tensor in zip(args, tensor_mask, strict=True)
+        if is_tensor
+    )
+    if total_bytes * pool_size > max_clone_bytes:
+        return None
+
+    pool: list[tuple[object, ...]] = []
+    for _ in range(pool_size):
+        pool.append(
+            tuple(arg.clone() if is_tensor else arg for arg, is_tensor in zip(args, tensor_mask, strict=True))
+        )
+    return pool
+
+
+def cuda_time_ms(
+    fn: Callable[..., object],
+    *,
+    args: tuple[object, ...] = (),
+    warmup: int = 10,
+    iterations: int = 50,
+    trials: int = 3,
+    flush_l2: bool = True,
+    clone_pool: bool = True,
+    clone_pool_size: int = 3,
+    clone_pool_max_bytes: int = 1 << 30,
+) -> float:
+    cache = _get_l2_flush_cache() if flush_l2 else None
+    arg_pool = _build_arg_pool(args, pool_size=clone_pool_size, max_clone_bytes=clone_pool_max_bytes) if clone_pool else None
+
+    def run_once() -> None:
+        if cache is not None:
+            cache.zero_()
+        if arg_pool is None:
+            fn(*args)
+        else:
+            current_args = arg_pool[run_once.iteration % len(arg_pool)]
+            fn(*current_args)
+
+    run_once.iteration = 0  # type: ignore[attr-defined]
+
+    for _ in range(warmup):
+        run_once()
+    torch.cuda.synchronize()
+
+    trial_means: list[float] = []
+
+    def on_trace_ready(prof) -> None:
+        total_ms = _sum_cuda_kernel_ms(prof.profiler.kineto_results)
+        trial_means.append(total_ms / iterations)
+
+    try:
+        schedule = torch.profiler.schedule(wait=0, warmup=1, active=1, repeat=trials)
+        with torch.profiler.profile(
+            activities=[torch.profiler.ProfilerActivity.CUDA],
+            schedule=schedule,
+            on_trace_ready=on_trace_ready,
+            acc_events=True,
+        ) as prof:
+            for _ in range(trials):
+                for _ in range(iterations):
+                    run_once.iteration += 1  # type: ignore[attr-defined]
+                    run_once()
+                prof.step()
+                for _ in range(iterations):
+                    run_once.iteration += 1  # type: ignore[attr-defined]
+                    run_once()
+                prof.step()
+    except (AttributeError, RuntimeError):
+        pass
+
+    if trial_means and not any(total_ms > 0 for total_ms in trial_means):
+        trial_means.clear()
+
+    if not trial_means:
+        for _ in range(trials):
+            start_events = [torch.cuda.Event(enable_timing=True) for _ in range(iterations)]
+            end_events = [torch.cuda.Event(enable_timing=True) for _ in range(iterations)]
+            for idx in range(iterations):
+                if cache is not None:
+                    cache.zero_()
+                start_events[idx].record()
+                if arg_pool is None:
+                    fn(*args)
+                else:
+                    fn(*arg_pool[idx % len(arg_pool)])
+                end_events[idx].record()
+            torch.cuda.synchronize()
+            times = [start.elapsed_time(end) for start, end in zip(start_events, end_events, strict=True)]
+            trial_means.append(sum(times) / len(times))
+
+    trial_means.sort()
+    return trial_means[len(trial_means) // 2]
--- a/python/operator_runtime_testing/profiler.py
+++ b/python/operator_runtime_testing/profiler.py
@ -9,23 +9,9 @@ class PerformanceResult:
    backend: str
    shape: str
    dtype: str
-    bytes: int
-    flops: int
    runtime_ms: float
    torch_ms: float | None = None

-    @property
-    def gbytes_per_sec(self) -> float:
-        if self.runtime_ms <= 0:
-            return 0.0
-        return self.bytes / (1024**3) / (self.runtime_ms / 1000.0)
-
-    @property
-    def gflops_per_sec(self) -> float:
-        if self.runtime_ms <= 0:
-            return 0.0
-        return self.flops / 1e9 / (self.runtime_ms / 1000.0)
-
    @property
    def speedup(self) -> float | None:
        if self.torch_ms is None or self.runtime_ms <= 0:
--- a/scripts/build_metax.sh
+++ b/scripts/build_metax.sh
@ -0,0 +1,72 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+BUILD_DIR="${BUILD_DIR:-${ROOT}/build-metax}"
+MACA_PATH="${MACA_PATH:-/opt/maca}"
+CUCC_PATH="${CUCC_PATH:-${MACA_PATH}/tools/cu-bridge}"
+CUDA_PATH="${CUDA_PATH:-${CUCC_PATH}}"
+PYTHON_BIN="${PYTHON_BIN:-python3}"
+
+MODE="${1:-build}"
+
+export MACA_PATH
+export CUCC_PATH
+export CUDA_PATH
+export LD_LIBRARY_PATH="${MACA_PATH}/lib:${LD_LIBRARY_PATH:-}"
+
+configure() {
+  mkdir -p "${BUILD_DIR}"
+  "${PYTHON_BIN}" -m cmake -G Ninja "${ROOT}" \
+    -B "${BUILD_DIR}" \
+    -DCAMP_ENABLE_NVIDIA=OFF \
+    -DCAMP_ENABLE_METAX=ON
+}
+
+build() {
+  "${PYTHON_BIN}" -m cmake --build "${BUILD_DIR}" -- -v
+}
+
+test_pytest() {
+  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" -m pytest "${ROOT}/tests" -v --backend metax
+}
+
+test_run_ops() {
+  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/tests/run_ops.py" --op all --backend metax --mode all
+}
+
+test_examples() {
+  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/01_copy.py" --backend metax
+  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/02_vector_add.py" --backend metax
+  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/03_reduce_sum.py" --backend metax
+  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/04_softmax.py" --backend metax
+}
+
+case "${MODE}" in
+  configure)
+    configure
+    ;;
+  build)
+    configure
+    build
+    ;;
+  test)
+    test_pytest
+    test_run_ops
+    test_examples
+    ;;
+  all)
+    configure
+    build
+    test_pytest
+    test_run_ops
+    test_examples
+    ;;
+  clean)
+    rm -rf "${BUILD_DIR}"
+    ;;
+  *)
+    echo "usage: $0 {configure|build|test|all|clean}" >&2
+    exit 2
+    ;;
+esac
--- a/tests/bench/copy.py
+++ b/tests/bench/copy.py
@ -12,32 +12,25 @@ if str(ROOT) not in sys.path:

 import torch

-from operator_runtime import copy
-from operator_runtime.testing import cuda_time_ms, PerformanceResult
+from operator_runtime_testing import cuda_time_ms, PerformanceResult
 from tests.cases import copy as copy_cases


-def _estimate_copy(tensor: torch.Tensor) -> tuple[int, int]:
-    elem_bytes = tensor.element_size()
-    return 2 * tensor.numel() * elem_bytes, 0
-
-
 def bench_copy(backend: str) -> list[PerformanceResult]:
    rows: list[PerformanceResult] = []
    for case in copy_cases.benchmark_cases():
        src = torch.randn(case["shape"], dtype=case["dtype"], device="cuda")
        out = torch.empty_like(src)
-        runtime = cuda_time_ms(lambda: copy(src, backend=backend))
-        torch_ms = cuda_time_ms(lambda: out.copy_(src))
-        bytes_, flops = _estimate_copy(src)
+        from operator_runtime import prepare_copy
+        with prepare_copy(out, src, backend=backend) as prepared:
+            runtime = cuda_time_ms(prepared.run_inputs, args=(src,))
+        torch_ms = cuda_time_ms(lambda src: out.copy_(src), args=(src,))
        rows.append(
            PerformanceResult(
                "copy",
                backend,
                str(tuple(src.shape)),
                str(src.dtype),
-                bytes_,
-                flops,
                runtime,
                torch_ms,
            )
--- a/tests/bench/reduce_sum.py
+++ b/tests/bench/reduce_sum.py
@ -12,32 +12,25 @@ if str(ROOT) not in sys.path:

 import torch

-from operator_runtime import reduce_sum
-from operator_runtime.testing import cuda_time_ms, PerformanceResult
+from operator_runtime_testing import cuda_time_ms, PerformanceResult
 from tests.cases import reduce_sum as reduce_sum_cases


-def _estimate_reduce_sum(tensor: torch.Tensor) -> tuple[int, int]:
-    rows = tensor.shape[0]
-    elem_bytes = tensor.element_size()
-    return (tensor.numel() + rows) * elem_bytes, tensor.numel()
-
-
 def bench_reduce_sum(backend: str) -> list[PerformanceResult]:
    rows: list[PerformanceResult] = []
    for case in reduce_sum_cases.benchmark_cases():
        src = torch.randn(case["shape"], dtype=case["dtype"], device="cuda")
-        runtime = cuda_time_ms(lambda: reduce_sum(src, dim=1, backend=backend))
-        torch_ms = cuda_time_ms(lambda: torch.sum(src, dim=1))
-        bytes_, flops = _estimate_reduce_sum(src)
+        from operator_runtime import prepare_reduce_sum
+        out = torch.empty(src.shape[0], dtype=src.dtype, device="cuda")
+        with prepare_reduce_sum(out, src, dim=1, backend=backend) as prepared:
+            runtime = cuda_time_ms(prepared.run_inputs, args=(src,))
+        torch_ms = cuda_time_ms(lambda src: torch.sum(src, dim=1, out=out), args=(src,))
        rows.append(
            PerformanceResult(
                "reduce_sum",
                backend,
                str(tuple(src.shape)),
                str(src.dtype),
-                bytes_,
-                flops,
                runtime,
                torch_ms,
            )
--- a/tests/bench/softmax.py
+++ b/tests/bench/softmax.py
@ -12,31 +12,25 @@ if str(ROOT) not in sys.path:

 import torch

-from operator_runtime import softmax
-from operator_runtime.testing import cuda_time_ms, PerformanceResult
+from operator_runtime_testing import cuda_time_ms, PerformanceResult
 from tests.cases import softmax as softmax_cases


-def _estimate_softmax(tensor: torch.Tensor) -> tuple[int, int]:
-    elem_bytes = tensor.element_size()
-    return 5 * tensor.numel() * elem_bytes, 4 * tensor.numel()
-
-
 def bench_softmax(backend: str) -> list[PerformanceResult]:
    rows: list[PerformanceResult] = []
    for case in softmax_cases.benchmark_cases():
        src = torch.randn(case["shape"], dtype=case["dtype"], device="cuda")
-        runtime = cuda_time_ms(lambda: softmax(src, dim=1, backend=backend))
-        torch_ms = cuda_time_ms(lambda: torch.softmax(src, dim=1))
-        bytes_, flops = _estimate_softmax(src)
+        out = torch.empty_like(src)
+        from operator_runtime import prepare_softmax
+        with prepare_softmax(out, src, dim=1, backend=backend) as prepared:
+            runtime = cuda_time_ms(prepared.run_inputs, args=(src,))
+        torch_ms = cuda_time_ms(lambda src: torch.softmax(src, dim=1, out=out), args=(src,))
        rows.append(
            PerformanceResult(
                "softmax",
                backend,
                str(tuple(src.shape)),
                str(src.dtype),
-                bytes_,
-                flops,
                runtime,
                torch_ms,
            )
--- a/tests/bench/vector_add.py
+++ b/tests/bench/vector_add.py
@ -12,32 +12,26 @@ if str(ROOT) not in sys.path:

 import torch

-from operator_runtime import vector_add
-from operator_runtime.testing import cuda_time_ms, PerformanceResult
+from operator_runtime_testing import cuda_time_ms, PerformanceResult
 from tests.cases import vector_add as vector_add_cases


-def _estimate_vector_add(tensor: torch.Tensor) -> tuple[int, int]:
-    elem_bytes = tensor.element_size()
-    return 3 * tensor.numel() * elem_bytes, tensor.numel()
-
-
 def bench_vector_add(backend: str) -> list[PerformanceResult]:
    rows: list[PerformanceResult] = []
    for case in vector_add_cases.benchmark_cases():
        a = torch.randn(case["shape"], dtype=case["dtype"], device="cuda")
        b = torch.randn_like(a)
-        runtime = cuda_time_ms(lambda: vector_add(a, b, backend=backend))
-        torch_ms = cuda_time_ms(lambda: torch.add(a, b))
-        bytes_, flops = _estimate_vector_add(a)
+        out = torch.empty_like(a)
+        from operator_runtime import prepare_vector_add
+        with prepare_vector_add(out, a, b, backend=backend) as prepared:
+            runtime = cuda_time_ms(prepared.run_inputs, args=(a, b))
+        torch_ms = cuda_time_ms(lambda a, b: torch.add(a, b, out=out), args=(a, b))
        rows.append(
            PerformanceResult(
                "vector_add",
                backend,
                str(tuple(a.shape)),
                str(a.dtype),
-                bytes_,
-                flops,
                runtime,
                torch_ms,
            )
--- a/tests/cases/copy.py
+++ b/tests/cases/copy.py
@ -22,5 +22,10 @@ def api_error_cases():

 def benchmark_cases():
    return [
-        {"name": "contiguous_1m", "shape": (1 << 20,), "dtype": torch.float16},
+        {"name": "contiguous_1001k_fp16", "shape": (1001 * 1024,), "dtype": torch.float16},
+        {"name": "contiguous_4093k_fp16", "shape": (4093 * 1024,), "dtype": torch.float16},
+        {"name": "contiguous_65521k_fp16", "shape": (65521 * 1024,), "dtype": torch.float16},
+        {"name": "contiguous_1001k_fp32", "shape": (1001 * 1024,), "dtype": torch.float32},
+        {"name": "contiguous_4093k_fp32", "shape": (4093 * 1024,), "dtype": torch.float32},
+        {"name": "contiguous_65521k_fp32", "shape": (65521 * 1024,), "dtype": torch.float32},
    ]
--- a/tests/cases/reduce_sum.py
+++ b/tests/cases/reduce_sum.py
@ -24,5 +24,10 @@ def api_error_cases():

 def benchmark_cases():
    return [
-        {"name": "rowwise_1024x1024", "shape": (1024, 1024), "dtype": torch.float32},
+        {"name": "rowwise_144x1280", "shape": (144, 1280), "dtype": torch.float32},
+        {"name": "rowwise_1008x1280", "shape": (1008, 1280), "dtype": torch.float32},
+        {"name": "rowwise_1008x2304", "shape": (1008, 2304), "dtype": torch.float32},
+        {"name": "rowwise_2032x2304", "shape": (2032, 2304), "dtype": torch.float32},
+        {"name": "rowwise_784x3840", "shape": (784, 3840), "dtype": torch.float32},
+        {"name": "rowwise_4080x1536", "shape": (4080, 1536), "dtype": torch.float32},
    ]
--- a/tests/cases/softmax.py
+++ b/tests/cases/softmax.py
@ -24,5 +24,10 @@ def api_error_cases():

 def benchmark_cases():
    return [
-        {"name": "rowwise_1024x1024", "shape": (1024, 1024), "dtype": torch.float32},
+        {"name": "rowwise_144x1280", "shape": (144, 1280), "dtype": torch.float32},
+        {"name": "rowwise_1008x1280", "shape": (1008, 1280), "dtype": torch.float32},
+        {"name": "rowwise_1008x2304", "shape": (1008, 2304), "dtype": torch.float32},
+        {"name": "rowwise_2032x2304", "shape": (2032, 2304), "dtype": torch.float32},
+        {"name": "rowwise_784x3840", "shape": (784, 3840), "dtype": torch.float32},
+        {"name": "rowwise_4080x1536", "shape": (4080, 1536), "dtype": torch.float32},
    ]
--- a/tests/cases/vector_add.py
+++ b/tests/cases/vector_add.py
@ -22,5 +22,10 @@ def api_error_cases():

 def benchmark_cases():
    return [
-        {"name": "contiguous_1m", "shape": (1 << 20,), "dtype": torch.float16},
+        {"name": "contiguous_1001k_fp16", "shape": (1001 * 1024,), "dtype": torch.float16},
+        {"name": "contiguous_4093k_fp16", "shape": (4093 * 1024,), "dtype": torch.float16},
+        {"name": "contiguous_65521k_fp16", "shape": (65521 * 1024,), "dtype": torch.float16},
+        {"name": "contiguous_1001k_fp32", "shape": (1001 * 1024,), "dtype": torch.float32},
+        {"name": "contiguous_4093k_fp32", "shape": (4093 * 1024,), "dtype": torch.float32},
+        {"name": "contiguous_65521k_fp32", "shape": (65521 * 1024,), "dtype": torch.float32},
    ]
--- a/tests/op_tests/init.py
+++ b/tests/op_tests/init.py
--- a/tests/op_tests/test_copy.py
+++ b/tests/op_tests/test_copy.py
@ -4,7 +4,7 @@ import pytest
 import torch

 from operator_runtime import copy, copy_
-from operator_runtime.testing import assert_close, require_cuda
+from operator_runtime_testing import assert_close, require_cuda
 from tests.cases import copy as copy_cases


--- a/tests/op_tests/test_reduce_sum.py
+++ b/tests/op_tests/test_reduce_sum.py
@ -4,7 +4,7 @@ import pytest
 import torch

 from operator_runtime import reduce_sum, reduce_sum_
-from operator_runtime.testing import assert_close, require_cuda
+from operator_runtime_testing import assert_close, require_cuda
 from tests.cases import reduce_sum as reduce_sum_cases


--- a/tests/op_tests/test_softmax.py
+++ b/tests/op_tests/test_softmax.py
@ -4,7 +4,7 @@ import pytest
 import torch

 from operator_runtime import softmax, softmax_
-from operator_runtime.testing import assert_close, require_cuda
+from operator_runtime_testing import assert_close, require_cuda
 from tests.cases import softmax as softmax_cases


--- a/tests/op_tests/test_vector_add.py
+++ b/tests/op_tests/test_vector_add.py
@ -5,7 +5,7 @@ import torch

 from operator_runtime import vector_add, vector_add_
 from operator_runtime.ops.vector_add import prepare_vector_add
-from operator_runtime.testing import assert_close, require_cuda
+from operator_runtime_testing import assert_close, require_cuda
 from tests.cases import vector_add as vector_add_cases


@ -45,7 +45,7 @@ def test_vector_add_api_contract(case, backend):

 def test_prepared_vector_add_reuses_descriptor(backend):
    require_cuda()
-    if backend != "nvidia":
+    if backend == "tilelang":
        pytest.skip("descriptor lifecycle test targets C ABI backend")
    a = torch.randn((1024,), device="cuda")
    b = torch.randn((1024,), device="cuda")
--- a/tests/run_ops.py
+++ b/tests/run_ops.py
@ -16,7 +16,7 @@ if str(ROOT) not in sys.path:


 def _discover_ops() -> tuple[str, ...]:
-    ops_dir = ROOT / "tests" / "ops"
+    ops_dir = ROOT / "tests" / "op_tests"
    bench_dir = ROOT / "tests" / "bench"
    names: set[str] = set()

@ -58,8 +58,6 @@ def _format_bench_table(rows) -> str:
        "runtime_ms",
        "torch_ms",
        "speedup",
-        "GB/s",
-        "GFLOP/s",
    ]
    body = []
    for row in rows:
@ -74,8 +72,6 @@ def _format_bench_table(rows) -> str:
                f"{row.runtime_ms:.4f}",
                torch_ms,
                speedup,
-                f"{row.gbytes_per_sec:.2f}",
-                f"{row.gflops_per_sec:.2f}",
            ]
        )

@ -114,7 +110,7 @@ def main() -> int:
    ops = _discover_ops()
    parser = argparse.ArgumentParser()
    parser.add_argument("--op", choices=[*ops, "all"], default="all")
-    parser.add_argument("--backend", choices=["nvidia", "tilelang"], default="nvidia")
+    parser.add_argument("--backend", choices=["nvidia", "tilelang", "metax"], default="nvidia")
    parser.add_argument("--mode", choices=["test", "bench", "all"], default="all")
    args = parser.parse_args()

@ -129,7 +125,7 @@ def main() -> int:

    for op in selected_ops:
        if args.mode in ("test", "all"):
-            ok, detail = _run_pytest(f"tests/ops/test_{op}.py", args.backend)
+            ok, detail = _run_pytest(f"tests/op_tests/test_{op}.py", args.backend)
            rows.append(["test", op, args.backend, "ok" if ok else "fail", detail])
            failed = failed or not ok
        if args.mode in ("bench", "all"):
--- a/tests/test_tilelang_copy_templates.py
+++ b/tests/test_tilelang_copy_templates.py
@ -1,58 +0,0 @@
-from __future__ import annotations
-
-import importlib.util
-
-import pytest
-import torch
-
-from operator_runtime.testing import require_cuda
-
-pytestmark = pytest.mark.skipif(
-    importlib.util.find_spec("tilelang") is None,
-    reason="tilelang is not installed",
-)
-
-
-@pytest.mark.parametrize(
-    ("copy_fn_name", "copy_out_fn_name", "module_name"),
-    [
-        ("copy_eager", "copy_eager_", "ops.common.tilelang.eager_copy"),
-        ("copy_lazy_out_idx", "copy_lazy_out_idx_", "ops.common.tilelang.lazy_out_idx_copy"),
-    ],
-)
-def test_tilelang_copy_templates_match_torch(copy_fn_name, copy_out_fn_name, module_name) -> None:
-    require_cuda()
-    module = __import__(module_name, fromlist=[copy_fn_name, copy_out_fn_name])
-    copy_fn = getattr(module, copy_fn_name)
-    copy_out_fn = getattr(module, copy_out_fn_name)
-
-    src = torch.randn((1024,), dtype=torch.float32, device="cuda")
-    out = copy_fn(src)
-    torch.testing.assert_close(out, src)
-
-    user_out = torch.empty_like(src)
-    returned = copy_out_fn(user_out, src)
-    assert returned is user_out
-    torch.testing.assert_close(user_out, src)
-
-
-@pytest.mark.parametrize(
-    ("copy_out_fn_name", "module_name", "message"),
-    [
-        ("copy_eager_", "ops.common.tilelang.eager_copy", "matching shapes"),
-        ("copy_lazy_out_idx_", "ops.common.tilelang.lazy_out_idx_copy", "matching shapes"),
-    ],
-)
-def test_tilelang_copy_templates_reject_mismatched_out(
-    copy_out_fn_name,
-    module_name,
-    message,
-) -> None:
-    require_cuda()
-    module = __import__(module_name, fromlist=[copy_out_fn_name])
-    copy_out_fn = getattr(module, copy_out_fn_name)
-
-    src = torch.randn((1024,), dtype=torch.float32, device="cuda")
-    out = torch.empty((512,), dtype=torch.float32, device="cuda")
-    with pytest.raises(ValueError, match=message):
-        copy_out_fn(out, src)
Author	SHA1	Message	Date
yutianyu	757dbbec74	refactor: split operator headers into public and detail layers	2026-05-08 11:18:47 +08:00
yutianyu	6e02c3eb55	feat: add MACA-based MetaX backend support Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-08 09:49:13 +08:00
yutianyu	ec8ed33b27	docs: add detailed info md Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	6d1f205610	refactor: scaffold phase1 kernel skeletons Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	7d94bb17f9	docs: add phase1 kernel writing guide Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	9147a4821d	feat: add kernel-focused benchmark timing harness Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	346fd7db12	refactor: rename op tests and update benchmark entrypoint Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	a2cda3fda7	feat: add more shared test cases Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	e9a4c87564	docs: update guides for refactored python layout Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00
yutianyu	dcf4d25896	refactor: split runtime internals and testing package Co-authored-by: wawahejun <hejunlbbc@gmail.com>	2026-05-05 16:50:18 +08:00