108 changed files with 1 additions and 6323 deletions
--- a/.cnb.yml
+++ b/.cnb.yml
@ -1,9 +0,0 @@
-$:
-  vscode:
-    - docker:
-        image: docker.cnb.cool/yutianyu.yi/image/ncu
-      runner:
-        tags: cnb:arch:amd64:gpu
-      services:
-        - vscode
-        - docker
--- a/.gitignore
+++ b/.gitignore
@ -1,14 +0,0 @@
-build/
-build-*/
-*.egg-info/
-__pycache__/
-.pytest_cache/
-.mypy_cache/
-.ruff_cache/
-*.pyc
-*.so
-*.o
-*.a
-*.log
-/generated/
-/third_party/
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@ -1,114 +0,0 @@
-cmake_minimum_required(VERSION 3.22)
-
-project(operator_runtime_training LANGUAGES CXX)
-
-option(CAMP_ENABLE_NVIDIA "Build NVIDIA CUDA backend" ON)
-option(CAMP_ENABLE_TILELANG "Enable TileLang Python backend metadata" ON)
-option(CAMP_ENABLE_METAX "Build MetaX backend" OFF)
-set(CAMP_ENABLE_CUTE "AUTO" CACHE STRING "Enable optional CuTe/CUTLASS headers for NVIDIA custom operators: AUTO, ON, or OFF")
-set_property(CACHE CAMP_ENABLE_CUTE PROPERTY STRINGS AUTO ON OFF)
-
-set(CAMP_CUTLASS_ROOT "" CACHE PATH "Optional CUTLASS checkout root for CuTe-based custom NVIDIA operators")
-set(CAMP_CUTE_INCLUDE_DIRS "" CACHE STRING "Optional semicolon-separated CuTe/CUTLASS include directories")
-
-set(CMAKE_CXX_STANDARD 17)
-set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CUDA_STANDARD 17)
-set(CMAKE_CUDA_STANDARD_REQUIRED ON)
-set(CMAKE_POSITION_INDEPENDENT_CODE ON)
-
-list(PREPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
-
-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()
-
-string(TOUPPER "${CAMP_ENABLE_CUTE}" CAMP_ENABLE_CUTE_MODE)
-if(NOT CAMP_ENABLE_CUTE_MODE MATCHES "^(AUTO|ON|OFF)$")
-  message(FATAL_ERROR "CAMP_ENABLE_CUTE must be AUTO, ON, or OFF")
-endif()
-
-function(camp_append_cute_root root)
-  if(NOT root)
-    return()
-  endif()
-  list(APPEND CAMP_CUTE_CANDIDATE_INCLUDE_DIRS
-    "${root}"
-    "${root}/include"
-    "${root}/tools/util/include"
-  )
-  set(CAMP_CUTE_CANDIDATE_INCLUDE_DIRS "${CAMP_CUTE_CANDIDATE_INCLUDE_DIRS}" PARENT_SCOPE)
-endfunction()
-
-set(CAMP_CUTE_TARGET_INCLUDE_DIRS "")
-set(CAMP_USE_CUTE OFF)
-if(NOT CAMP_ENABLE_CUTE_MODE STREQUAL "OFF")
-  if(NOT CAMP_ENABLE_NVIDIA)
-    if(CAMP_ENABLE_CUTE_MODE STREQUAL "ON")
-      message(FATAL_ERROR "CAMP_ENABLE_CUTE=ON requires CAMP_ENABLE_NVIDIA=ON")
-    endif()
-  else()
-    set(CAMP_CUTE_CANDIDATE_INCLUDE_DIRS "")
-
-    camp_append_cute_root("${CAMP_CUTLASS_ROOT}")
-    camp_append_cute_root("${CMAKE_CURRENT_SOURCE_DIR}/third_party/cutlass")
-    foreach(_camp_cute_env_var CUTLASS_ROOT CUTLASS_HOME CUTLASS_PATH)
-      if(DEFINED ENV{${_camp_cute_env_var}} AND NOT "$ENV{${_camp_cute_env_var}}" STREQUAL "")
-        camp_append_cute_root("$ENV{${_camp_cute_env_var}}")
-      endif()
-    endforeach()
-
-    if(CAMP_CUTE_INCLUDE_DIRS)
-      list(APPEND CAMP_CUTE_CANDIDATE_INCLUDE_DIRS ${CAMP_CUTE_INCLUDE_DIRS})
-    endif()
-
-    if(CAMP_CUTE_CANDIDATE_INCLUDE_DIRS)
-      list(REMOVE_DUPLICATES CAMP_CUTE_CANDIDATE_INCLUDE_DIRS)
-    endif()
-
-    foreach(_camp_cute_include_dir IN LISTS CAMP_CUTE_CANDIDATE_INCLUDE_DIRS)
-      if(EXISTS "${_camp_cute_include_dir}/cute/tensor.hpp")
-        list(APPEND CAMP_CUTE_TARGET_INCLUDE_DIRS "${_camp_cute_include_dir}")
-      endif()
-    endforeach()
-
-    if(CAMP_CUTE_TARGET_INCLUDE_DIRS)
-      list(REMOVE_DUPLICATES CAMP_CUTE_TARGET_INCLUDE_DIRS)
-      set(CAMP_USE_CUTE ON)
-      message(STATUS "CuTe/CUTLASS support enabled: ${CAMP_CUTE_TARGET_INCLUDE_DIRS}")
-    elseif(CAMP_ENABLE_CUTE_MODE STREQUAL "ON")
-      message(FATAL_ERROR "CAMP_ENABLE_CUTE=ON requires cute/tensor.hpp via CAMP_CUTLASS_ROOT, CAMP_CUTE_INCLUDE_DIRS, or CUTLASS_ROOT/CUTLASS_HOME/CUTLASS_PATH")
-    else()
-      message(STATUS "CuTe/CUTLASS support disabled: cute/tensor.hpp was not found")
-    endif()
-  endif()
-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()
-endif()
-
-if(CAMP_ENABLE_METAX)
-  set(CAMP_METAX_COMPILE_OPTIONS
-    -D__MACA_NO_HALF_OPERATORS__
-    -D__FAST_BLOCK_RED__
-  )
-endif()
-
-add_subdirectory(ops)
--- a/CMakePresets.json
+++ b/CMakePresets.json
@ -1,76 +0,0 @@
-{
-  "version": 6,
-  "cmakeMinimumRequired": {
-    "major": 3,
-    "minor": 22,
-    "patch": 0
-  },
-  "configurePresets": [
-    {
-      "name": "base",
-      "hidden": true,
-      "generator": "Ninja",
-      "binaryDir": "${sourceDir}/build/${presetName}",
-      "cacheVariables": {
-        "CMAKE_CXX_STANDARD": "17",
-        "CMAKE_CUDA_STANDARD": "17"
-      }
-    },
-    {
-      "name": "nvidia-debug",
-      "inherits": "base",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Debug",
-        "CAMP_ENABLE_NVIDIA": "ON",
-        "CAMP_ENABLE_METAX": "OFF",
-        "CAMP_ENABLE_CUTE": "AUTO",
-        "CMAKE_CUDA_ARCHITECTURES": "native"
-      }
-    },
-    {
-      "name": "nvidia-release",
-      "inherits": "base",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "CAMP_ENABLE_NVIDIA": "ON",
-        "CAMP_ENABLE_METAX": "OFF",
-        "CAMP_ENABLE_CUTE": "AUTO",
-        "CMAKE_CUDA_ARCHITECTURES": "native"
-      }
-    },
-    {
-      "name": "nvidia-cute",
-      "inherits": "nvidia-release",
-      "cacheVariables": {
-        "CAMP_ENABLE_CUTE": "ON"
-      }
-    },
-    {
-      "name": "metax-release",
-      "inherits": "base",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "CAMP_ENABLE_NVIDIA": "OFF",
-        "CAMP_ENABLE_METAX": "ON"
-      }
-    }
-  ],
-  "buildPresets": [
-    {
-      "name": "nvidia-debug",
-      "configurePreset": "nvidia-debug"
-    },
-    {
-      "name": "nvidia-release",
-      "configurePreset": "nvidia-release"
-    },
-    {
-      "name": "nvidia-cute",
-      "configurePreset": "nvidia-cute"
-    },
-    {
-      "name": "metax-release",
-      "configurePreset": "metax-release"
-    }
-  ]
-}
--- a/README.md
+++ b/README.md
@ -1,175 +1,2 @@
-# Operator Runtime Training Camp
+# Intro-ops

-A nano GPU operator runtime for learning kernel development. The framework uses a backend-agnostic public C API with separate NVIDIA and MetaX build variants, while Python, tests, and benchmarks sit on top of the same runtime contract.
-
-## Quick Start
-
-```bash
-# activate environment
-conda activate py312
-
-# build NVIDIA (auto-fetches CUTLASS on first run)
-bash scripts/build_nvidia.sh build
-
-# run all NVIDIA tests
-bash scripts/build_nvidia.sh test
-
-# build MetaX
-bash scripts/build_metax.sh build
-
-# run MetaX tests
-bash scripts/build_metax.sh test
-
-# clean build
-bash scripts/build_nvidia.sh clean
-bash scripts/build_metax.sh clean
-```
-
-## Common Commands
-
-### Build Script Modes
-
-```bash
-bash scripts/build_nvidia.sh env          # show current build environment variables
-bash scripts/build_nvidia.sh configure    # cmake configure only
-bash scripts/build_nvidia.sh build        # configure + build
-bash scripts/build_nvidia.sh test         # run pytest + run_ops + examples
-bash scripts/build_nvidia.sh all          # build + test
-bash scripts/build_nvidia.sh clean        # remove build directory
-
-bash scripts/build_metax.sh env           # show current MACA build environment
-bash scripts/build_metax.sh configure     # cmake configure only
-bash scripts/build_metax.sh build         # configure + build
-bash scripts/build_metax.sh test          # run pytest + run_ops + examples
-bash scripts/build_metax.sh all           # build + test
-bash scripts/build_metax.sh clean         # remove build directory
-```
-
-### Single Operator Testing
-
-```bash
-# correctness test for one operator
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v
-
-# run_ops supports --mode: test, bench, all
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode test
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode bench
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode all
-
-# all operators
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op all --backend nvidia --mode all
-
-# MetaX
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-metax pytest tests/op_tests/test_copy.py -v --backend metax
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-metax python tests/run_ops.py --op all --backend metax --mode all
-```
-
-### Force Rebuild
-
-```bash
-# clear cmake cache and rebuild
-CAMP_FORCE_RECONFIGURE=1 bash scripts/build_nvidia.sh build
-
-# target specific GPU architecture
-CMAKE_CUDA_ARCHITECTURES=89 bash scripts/build_nvidia.sh build
-```
-
-## Project Structure
-
-```
-include/operator_runtime/
-  operator_runtime.h  <-- umbrella public runtime header
-  ops/<op>.h          <-- backend-agnostic public C API
-  detail/*.h          <-- shared internal helper layer
-ops/<op>/nvidia/
-  kernel.cuh          <-- implement your kernel here
-  <op>_cuda.cu        <-- unified public symbol implementation
-ops/<op>/metax/
-  <op>_metax.maca     <-- MetaX kernel + launch implementation
-ops/elementwise/<op>/
-  nvidia/*.cu         <-- elementwise NVIDIA implementations
-  metax/*.maca        <-- elementwise MetaX implementations
-python/operator_runtime/
-  ops/<op>.py         <-- Python wrapper over shared C API / TileLang
-tests/
-  op_tests/test_<op>.py   <-- correctness tests
-  cases/<op>.py           <-- test cases
-  bench/<op>.py           <-- benchmarks
-```
-
-## Architecture
-
- Public operator headers live under `include/operator_runtime/ops/*.h` and expose backend-agnostic symbols such as `oprt_create_copy_descriptor`.
- Backend-private code lives under `ops/.../<backend>/`.
- Python bindings use the same public symbol names for both compiled backends and select the active backend via `oprt_set_backend(...)` plus separate build outputs.
-
-## Operators
-
-| Operator | Difficulty | Key Concept |
-| --- | --- | --- |
-| `copy` | easy | grid-stride loop, vectorized memory access |
-| `vector_add` | easy | elementwise computation |
-| `reduce_sum` | medium | shared memory, warp reduction |
-| `softmax` | hard | multi-pass reduction + normalization |
-| `relu` | reference | elementwise framework (already implemented) |
-
-## Workflow
-
-1. Read the kernel scaffold in `ops/<op>/nvidia/kernel.cuh`
-2. Implement the TODO kernel
-3. Build: `bash scripts/build_nvidia.sh build`
-4. Test: `PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<op>.py -v`
-5. Benchmark: `PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op <op> --backend nvidia --mode bench`
-
-## Adding a New Operator
-
-See [docs/how-to-add-an-operator.md](docs/how-to-add-an-operator.md) for the full guide. Minimal steps:
-
-1. Create `ops/<op>/nvidia/` with kernel and cuda source
-2. Add Python bindings in `python/operator_runtime/ops/<op>.py`
-3. Add test cases, correctness tests, and benchmarks under `tests/`
-4. Re-run `bash scripts/build_nvidia.sh configure` when adding new NVIDIA `.cu` sources
-5. Build and verify
-
-## MetaX Backend
-
-The training runtime also supports a MetaX build variant through `bash scripts/build_metax.sh ...`.
-
-```bash
-# build MetaX variant
-bash scripts/build_metax.sh build
-
-# run MetaX correctness + benchmark flow
-bash scripts/build_metax.sh test
-```
-
-## TileLang on MetaX
-
-Stock pip `tilelang` is not sufficient for this machine. Use the source-built `/root/tilelang-metax` tree for TileLang-on-MetaX validation when needed.
-
-```bash
-# build /root/tilelang-metax separately with USE_MACA=ON
-
-CAMP_USE_TILELANG_METAX=1 \
-CAMP_TILELANG_SOURCE_ROOT=/root/tilelang-metax \
-bash scripts/build_metax.sh test
-```
-
-This mode exports:
- `PYTHONPATH=/root/tilelang-metax`
- `LD_LIBRARY_PATH=/root/tilelang-metax/build/lib:/opt/maca/lib:...`
-
-and validates `copy/vector_add/reduce_sum/softmax` with `backend=tilelang` on MetaX.
-
-## Environment Variables
-
-| Variable | Default | Description |
-| --- | --- | --- |
-| `CMAKE_CUDA_ARCHITECTURES` | `native` | Target GPU arch (e.g. `89` for L40) |
-| `CAMP_FORCE_RECONFIGURE` | `0` | Set `1` to clear CMake cache before configure |
-| `CAMP_ENABLE_CUTE` | `AUTO` | CuTe/CUTLASS support: `AUTO`, `ON`, `OFF` |
-| `BUILD_DIR` | `build-nvidia` | Build output directory |
-| `CAMP_BUILD_DIR` | — | Tell Python where to find `libcamp_ops.so` |
-| `CAMP_CUTLASS_ROOT` | — | Override CUTLASS path (skips auto-fetch) |
-| `CAMP_USE_TILELANG_METAX` | `0` | Set `1` to run TileLang-on-MetaX validation in `build_metax.sh` |
-| `CAMP_TILELANG_SOURCE_ROOT` | `/root/tilelang-metax` | Source-built TileLang tree used for MetaX TileLang validation |
--- a/README.zh.md
+++ b/README.zh.md
@ -1,175 +0,0 @@
-# Intro-ops 训练营
-
-面向 kernel 开发的 nano 级 GPU 算子运行时。框架现在采用后端无关的公共 C API，并提供独立的 NVIDIA 与 MetaX 构建变体；Python、测试和 benchmark 都建立在同一套运行时契约之上。
-
-## 快速开始
-
-```bash
-# 激活环境
-conda activate py312
-
-# 构建 NVIDIA（首次运行自动拉取 CUTLASS）
-bash scripts/build_nvidia.sh build
-
-# 运行全部 NVIDIA 测试
-bash scripts/build_nvidia.sh test
-
-# 构建 MetaX
-bash scripts/build_metax.sh build
-
-# 运行 MetaX 测试
-bash scripts/build_metax.sh test
-
-# 清理构建
-bash scripts/build_nvidia.sh clean
-bash scripts/build_metax.sh clean
-```
-
-## 常用命令
-
-### 构建脚本模式
-
-```bash
-bash scripts/build_nvidia.sh env          # 显示当前构建环境变量
-bash scripts/build_nvidia.sh configure    # 仅 cmake 配置
-bash scripts/build_nvidia.sh build        # 配置 + 编译
-bash scripts/build_nvidia.sh test         # 运行 pytest + run_ops + examples
-bash scripts/build_nvidia.sh all          # 编译 + 测试
-bash scripts/build_nvidia.sh clean        # 删除构建目录
-
-bash scripts/build_metax.sh env           # 显示当前 MACA 构建环境
-bash scripts/build_metax.sh configure     # 仅 cmake 配置
-bash scripts/build_metax.sh build         # 配置 + 编译
-bash scripts/build_metax.sh test          # 运行 pytest + run_ops + examples
-bash scripts/build_metax.sh all           # 编译 + 测试
-bash scripts/build_metax.sh clean         # 删除构建目录
-```
-
-### 单算子测试
-
-```bash
-# 单算子正确性测试
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v
-
-# run_ops 支持 --mode: test, bench, all
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode test
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode bench
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode all
-
-# 全部算子
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op all --backend nvidia --mode all
-
-# MetaX
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-metax pytest tests/op_tests/test_copy.py -v --backend metax
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-metax python tests/run_ops.py --op all --backend metax --mode all
-```
-
-### 强制重新构建
-
-```bash
-# 清除 cmake 缓存并重新构建
-CAMP_FORCE_RECONFIGURE=1 bash scripts/build_nvidia.sh build
-
-# 指定 GPU 架构
-CMAKE_CUDA_ARCHITECTURES=89 bash scripts/build_nvidia.sh build
-```
-
-## 项目结构
-
-```
-include/operator_runtime/
-  operator_runtime.h  <-- 公共运行时总头
-  ops/<op>.h          <-- 后端无关的公开 C API
-  detail/*.h          <-- 共享内部 helper 分层
-ops/<op>/nvidia/
-  kernel.cuh          <-- 在这里实现你的 kernel
-  <op>_cuda.cu        <-- 统一公共符号实现
-ops/<op>/metax/
-  <op>_metax.maca     <-- MetaX kernel + launch 实现
-ops/elementwise/<op>/
-  nvidia/*.cu         <-- elementwise NVIDIA 实现
-  metax/*.maca        <-- elementwise MetaX 实现
-python/operator_runtime/
-  ops/<op>.py         <-- 基于共享 C API / TileLang 的 Python 包装
-tests/
-  op_tests/test_<op>.py   <-- 正确性测试
-  cases/<op>.py           <-- 测试用例
-  bench/<op>.py           <-- 性能基准
-```
-
-## 架构说明
-
- 公开算子头文件位于 `include/operator_runtime/ops/*.h`，暴露统一的后端无关符号，例如 `oprt_create_copy_descriptor`。
- 后端私有实现位于 `ops/.../<backend>/`。
- Python 绑定对 NVIDIA 和 MetaX 复用同一套公开符号，通过 `oprt_set_backend(...)` 和独立构建产物选择后端。
-
-## 算子列表
-
-| 算子 | 难度 | 核心概念 |
-| --- | --- | --- |
-| `copy` | 入门 | grid-stride loop、向量化内存访问 |
-| `vector_add` | 入门 | 逐元素计算 |
-| `reduce_sum` | 中等 | shared memory、warp 归约 |
-| `softmax` | 进阶 | 多趟归约 + 归一化 |
-| `relu` | 参考实现 | elementwise 框架（已实现） |
-
-## 开发流程
-
-1. 阅读 `ops/<op>/nvidia/kernel.cuh` 中的 kernel 骨架
-2. 实现 TODO kernel
-3. 构建：`bash scripts/build_nvidia.sh build`
-4. 测试：`PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<op>.py -v`
-5. Benchmark：`PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op <op> --backend nvidia --mode bench`
-
-## 新增算子
-
-完整指南见 [docs/how-to-add-an-operator.md](docs/how-to-add-an-operator.md)。最小步骤：
-
-1. 创建 `ops/<op>/nvidia/`，包含 kernel 和 cuda 源文件
-2. 在 `python/operator_runtime/ops/<op>.py` 添加 Python 绑定
-3. 在 `tests/` 下添加测试用例、正确性测试和 benchmark
-4. 新增 NVIDIA `.cu` 文件后，重新运行 `bash scripts/build_nvidia.sh configure`
-5. 构建并验证
-
-## MetaX 后端
-
-训练营同时支持 MetaX 构建变体，可通过 `bash scripts/build_metax.sh ...` 使用。
-
-```bash
-# 构建 MetaX 变体
-bash scripts/build_metax.sh build
-
-# 运行 MetaX 正确性与 benchmark 流程
-bash scripts/build_metax.sh test
-```
-
-## MetaX 上的 TileLang
-
-本机不能直接使用 pip 官方版 `tilelang`。只有在需要验证 MetaX 上的 TileLang 时，才使用源码构建的 `/root/tilelang-metax`。
-
-```bash
-# 先单独把 /root/tilelang-metax 以 USE_MACA=ON 构建完成
-
-CAMP_USE_TILELANG_METAX=1 \
-CAMP_TILELANG_SOURCE_ROOT=/root/tilelang-metax \
-bash scripts/build_metax.sh test
-```
-
-该模式会导出：
- `PYTHONPATH=/root/tilelang-metax`
- `LD_LIBRARY_PATH=/root/tilelang-metax/build/lib:/opt/maca/lib:...`
-
-并在 MetaX 上验证 `copy/vector_add/reduce_sum/softmax` 的 `backend=tilelang`。
-
-## 环境变量
-
-| 变量 | 默认值 | 说明 |
-| --- | --- | --- |
-| `CMAKE_CUDA_ARCHITECTURES` | `native` | 目标 GPU 架构（如 L40 用 `89`） |
-| `CAMP_FORCE_RECONFIGURE` | `0` | 设为 `1` 在配置前清除 CMake 缓存 |
-| `CAMP_ENABLE_CUTE` | `AUTO` | CuTe/CUTLASS 支持：`AUTO`、`ON`、`OFF` |
-| `BUILD_DIR` | `build-nvidia` | 构建输出目录 |
-| `CAMP_BUILD_DIR` | — | 告诉 Python 在哪里找 `libcamp_ops.so` |
-| `CAMP_CUTLASS_ROOT` | — | 覆盖 CUTLASS 路径（跳过自动拉取） |
-| `CAMP_USE_TILELANG_METAX` | `0` | 设为 `1` 时在 `build_metax.sh` 中执行 MetaX 上的 TileLang 验证 |
-| `CAMP_TILELANG_SOURCE_ROOT` | `/root/tilelang-metax` | 用于 MetaX TileLang 验证的源码版 TileLang 根目录 |
--- a/cmake/CMakeMACAInformation.cmake
+++ b/cmake/CMakeMACAInformation.cmake
@ -1,16 +0,0 @@
-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,82 +0,0 @@
-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/cmake/cuda_helpers.cmake
+++ b/cmake/cuda_helpers.cmake
@ -1,10 +0,0 @@
-if(NOT DEFINED CMAKE_CUDA_STANDARD)
-  set(CMAKE_CUDA_STANDARD 17)
-endif()
-
-if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-  set(CMAKE_CUDA_ARCHITECTURES native)
-endif()
-
-set(CMAKE_CUDA_STANDARD_REQUIRED ON)
-
--- a/course/README.md
+++ b/course/README.md
@ -1,18 +0,0 @@
-# 课程资料
-
-本目录用于存放训练营课件和学习记录。
-
-## 目录用途
-
- 课件、讲义、slides
- 个人学习笔记
- 实验记录和性能分析报告
-
-## 建议组织方式
-
-```
-course/
-  slides/           # 课件
-  notes/            # 学习笔记
-  reports/          # 实验报告
-```
--- a/docs/how-to-add-an-operator.md
+++ b/docs/how-to-add-an-operator.md
@ -1,131 +0,0 @@
-# 如何开发一个新算子
-
-## 目标
-
-新增一个算子的最小闭环包括：
-
-1. 在 `ops/<算子名>/nvidia/` 下实现 kernel 和 descriptor 生命周期。
-2. 在 `python/operator_runtime/ops/` 下补 Python API。
-3. 在 `tests/` 下补正确性测试和 benchmark。
-4. 重新构建并验证。
-
-训练营支持两条并行路径：
- **自定义算子**：`ops/<op>/nvidia/`，适合 kernel 教学、特殊 layout、reduce/softmax 等非逐元素结构。
- **elementwise 复用**：`ops/elementwise/<op>/nvidia/`，适合 unary/binary/broadcast 逐元素算子。
-
-## Step 1：明确算子接口
-
-开始前确认：
-
-1. 几个输入、几个输出
-2. 输出 shape 是否和输入一致
-3. 是否要求 dtype 一致
-4. 是否只支持 contiguous tensor
-5. 是否需要额外参数（`dim`、`scalar` 等）
-6. 是否需要 workspace
-
-## Step 2：创建算子目录
-
-最小文件集：
-
-```
-ops/<算子名>/nvidia/
-  kernel.cuh              # kernel 实现
-  <算子名>_cuda.cu        # descriptor 生命周期 + launch
-include/operator_runtime/ops/<算子名>.h   # C API 头文件
-python/operator_runtime/ops/<算子名>.py   # Python 绑定
-tests/cases/<算子名>.py                   # 测试数据
-tests/op_tests/test_<算子名>.py           # 正确性测试
-tests/bench/<算子名>.py                   # benchmark
-```
-
-### include 目录结构与创建原则
-
-```
-include/operator_runtime/
-  api.h               # 公共类型定义（status、dtype、tensor_view 等）
-  descriptor.h        # descriptor 基类
-  tensor_view.h       # tensor view 结构体
-  operator_runtime.h  # 汇总头文件
-  ops/
-    <算子名>.h        # 每个算子一个头文件，声明四个生命周期 C 函数
-  detail/
-    cuda_helpers.h    # CUDA 工具函数（blocks_for、stream 转换等）
-    elementwise.h     # elementwise descriptor helper
-    tensor_checks.h   # tensor 校验工具
-    operation.h       # operation 基类
-```
-
-创建原则：
-
- `ops/<算子名>.h` 是算子的公开 C API 契约，只声明四个生命周期函数（create/workspace/execute/destroy），不暴露实现细节。
- 所有函数使用 `extern "C"` + `OPRT_EXPORT`，保证 Python FFI 可以按符号名 dlsym。
- 参数类型只使用 `api.h` 中定义的公共类型（`oprt_status_t`、`oprt_tensor_view_t`、`oprt_operator_descriptor_t`、`oprt_stream_t`）。
- `detail/` 下放内部实现工具，不对外暴露，算子实现可以 include 但用户代码不应依赖。
- 新增算子只需在 `ops/` 下加一个头文件，不需要修改其他头文件。
-
-## Step 3：实现 NVIDIA 后端
-
-一个算子需要四个生命周期接口：
-
-1. `oprt_create_<op>_descriptor` — 检查输入合法性，保存运行信息
-2. `oprt_get_<op>_workspace_size` — 返回临时内存大小
-3. `oprt_execute_<op>` — 按 dtype dispatch 并 launch kernel
-4. `oprt_destroy_<op>_descriptor` — 释放 descriptor
-
-参考 `copy` 的实现：kernel 放 `kernel.cuh`，生命周期放 `<op>_cuda.cu`。
-
-如果是 elementwise 算子，复用 `ops/common/elementwise/nvidia/elementwise_nvidia.cuh` 和 `include/operator_runtime/detail/elementwise.h`，只需提供 device functor。参考 `relu`。
-
-## Step 4：补 Python 绑定
-
-Python 入口放在 `python/operator_runtime/ops/<算子名>.py`，通常暴露三个接口：
-
- `<op>` — out-of-place，自动分配输出
- `<op>_` — out-variant，调用方提供输出 tensor
- `prepare_<op>` — 创建 descriptor，支持多次执行复用
-
-新增后在 `ops/__init__.py` 和 `operator_runtime/__init__.py` 中导出。
-
-## Step 5：补测试
-
-`tests/cases/<算子名>.py` 组织测试数据，分三类：
-
- `correctness_cases()` — 正确性用例（shape、dtype、tolerance）
- `api_error_cases()` — 异常输入用例
- `benchmark_cases()` — 性能测试规模
-
-`tests/op_tests/test_<算子名>.py` 负责：
-
- 正确性对比（vs PyTorch）
- API contract 检查（shape/dtype 不匹配、非 contiguous）
- prepared 执行复用检查
-
-`tests/bench/<算子名>.py` 负责性能入口。
-
-## Step 6：构建和验证
-
-```bash
-# 新增 .cu 后必须重新 configure
-bash scripts/build_nvidia.sh configure
-bash scripts/build_nvidia.sh build
-
-# 单算子验证
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<算子名>.py -v --backend nvidia
-
-# 正确性 + benchmark
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op <算子名> --backend nvidia --mode all
-
-# 确认没有破坏已有算子
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op all --backend nvidia --mode test
-```
-
-## 现有模板参考
-
-| 模板 | 适用场景 |
-| --- | --- |
-| `copy` | 最简单的单输入单输出自定义算子 |
-| `vector_add` | 双输入 contiguous 自定义算子 |
-| `relu` | elementwise 框架复用，含标量参数 |
-| `reduce_sum` | 带 reduce 维度的算子 |
-| `softmax` | 多步归约 + 归一化 |
--- a/docs/phase1-kernel-writing.md
+++ b/docs/phase1-kernel-writing.md
@ -1,106 +0,0 @@
-# 第一阶段训练目标：写 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（需要先重新编译）
-bash scripts/build_nvidia.sh build
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<算子名>.py -v --backend nvidia
-
-# TileLang kernel（不需要编译）
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<算子名>.py -v --backend tilelang
-
-# 同时跑正确性 + benchmark
-PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op <算子名> --backend nvidia --mode all
-```
-
-四个算子两种后端全部通过，阶段一完成。
--- a/examples/01_copy.py
+++ b/examples/01_copy.py
@ -1,28 +0,0 @@
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[1]
-sys.path.insert(0, str(ROOT / "python"))
-sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime import copy
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
-    args = parser.parse_args()
-
-    src = torch.randn((1024,), device="cuda", dtype=torch.float32)
-    out = copy(src, backend=args.backend)
-    torch.testing.assert_close(out, src)
-    print(f"copy ok ({args.backend})")
-
-
-if __name__ == "__main__":
-    main()
--- a/examples/02_vector_add.py
+++ b/examples/02_vector_add.py
@ -1,29 +0,0 @@
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[1]
-sys.path.insert(0, str(ROOT / "python"))
-sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime import vector_add
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
-    args = parser.parse_args()
-
-    a = torch.randn((1024,), device="cuda", dtype=torch.float32)
-    b = torch.randn_like(a)
-    out = vector_add(a, b, backend=args.backend)
-    torch.testing.assert_close(out, a + b)
-    print(f"vector_add ok ({args.backend})")
-
-
-if __name__ == "__main__":
-    main()
--- a/examples/03_reduce_sum.py
+++ b/examples/03_reduce_sum.py
@ -1,28 +0,0 @@
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[1]
-sys.path.insert(0, str(ROOT / "python"))
-sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime import reduce_sum
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
-    args = parser.parse_args()
-
-    src = torch.randn((32, 128), device="cuda", dtype=torch.float32)
-    out = reduce_sum(src, dim=1, backend=args.backend)
-    torch.testing.assert_close(out, torch.sum(src, dim=1), atol=1e-5, rtol=1e-5)
-    print(f"reduce_sum ok ({args.backend})")
-
-
-if __name__ == "__main__":
-    main()
--- a/examples/04_softmax.py
+++ b/examples/04_softmax.py
@ -1,28 +0,0 @@
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[1]
-sys.path.insert(0, str(ROOT / "python"))
-sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime import softmax
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia", "tilelang", "metax"])
-    args = parser.parse_args()
-
-    src = torch.randn((32, 128), device="cuda", dtype=torch.float32)
-    out = softmax(src, dim=1, backend=args.backend)
-    torch.testing.assert_close(out, torch.softmax(src, dim=1), atol=1e-5, rtol=1e-5)
-    print(f"softmax ok ({args.backend})")
-
-
-if __name__ == "__main__":
-    main()
--- a/examples/05_tilelang_copy_modes.py
+++ b/examples/05_tilelang_copy_modes.py
@ -1,37 +0,0 @@
-from __future__ import annotations
-
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[1]
-sys.path.insert(0, str(ROOT / "python"))
-sys.path.insert(0, str(ROOT))
-
-import torch
-
-from ops.common.tilelang.eager_copy import copy_eager, copy_eager_
-from ops.common.tilelang.lazy_out_idx_copy import copy_lazy_out_idx, copy_lazy_out_idx_
-
-
-def main() -> None:
-    src = torch.randn((1024,), device="cuda", dtype=torch.float32)
-
-    eager = copy_eager(src)
-    torch.testing.assert_close(eager, src)
-
-    eager_out = torch.empty_like(src)
-    copy_eager_(eager_out, src)
-    torch.testing.assert_close(eager_out, src)
-
-    lazy = copy_lazy_out_idx(src)
-    torch.testing.assert_close(lazy, src)
-
-    lazy_out = torch.empty_like(src)
-    copy_lazy_out_idx_(lazy_out, src)
-    torch.testing.assert_close(lazy_out, src)
-
-    print("tilelang eager and lazy out_idx copy ok")
-
-
-if __name__ == "__main__":
-    main()
--- a/examples/06_relu.py
+++ b/examples/06_relu.py
@ -1,30 +0,0 @@
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[1]
-sys.path.insert(0, str(ROOT / "python"))
-sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime import relu
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser()
-    parser.add_argument("--backend", default="nvidia", choices=["nvidia"])
-    parser.add_argument("--negative-slope", type=float, default=0.0)
-    args = parser.parse_args()
-
-    src = torch.randn((1024,), device="cuda", dtype=torch.float32)
-    out = relu(src, negative_slope=args.negative_slope, backend=args.backend)
-    expected = torch.where(src > 0, src, src * args.negative_slope)
-    torch.testing.assert_close(out, expected)
-    print(f"relu ok ({args.backend}, negative_slope={args.negative_slope})")
-
-
-if __name__ == "__main__":
-    main()
--- a/include/operator_runtime/api.h
+++ b/include/operator_runtime/api.h
@ -1,56 +0,0 @@
-#pragma once
-
-#include <stdint.h>
-#include <stddef.h>
-
-#ifdef _WIN32
-#define OPRT_EXPORT __declspec(dllexport)
-#else
-#define OPRT_EXPORT __attribute__((visibility("default")))
-#endif
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-typedef enum {
-    OPRT_SUCCESS = 0,
-    OPRT_ERR_INVALID_ARG = 1,
-    OPRT_ERR_UNSUPPORTED_DTYPE = 2,
-    OPRT_ERR_RUNTIME = 3,
-    OPRT_ERR_INSUFFICIENT_WORKSPACE = 4,
-    OPRT_ERR_NOT_SUPPORTED = 5
-} oprt_status_t;
-
-typedef enum {
-    OPRT_DTYPE_F16 = 0,
-    OPRT_DTYPE_F32 = 1
-} oprt_dtype_t;
-
-typedef enum {
-    OPRT_BACKEND_AUTO = 0,
-    OPRT_BACKEND_NVIDIA = 1,
-    OPRT_BACKEND_METAX = 2
-} oprt_backend_t;
-
-#define OPRT_MAX_DIMS 8
-
-typedef void *oprt_stream_t;
-typedef struct oprt_operator_descriptor *oprt_operator_descriptor_t;
-
-typedef struct {
-    void *data;
-    oprt_dtype_t dtype;
-    int32_t ndim;
-    int64_t shape[OPRT_MAX_DIMS];
-    int64_t strides[OPRT_MAX_DIMS];
-} oprt_tensor_view_t;
-
-OPRT_EXPORT const char *oprt_status_string(oprt_status_t status);
-OPRT_EXPORT const char *oprt_backend_string(oprt_backend_t backend);
-OPRT_EXPORT oprt_status_t oprt_set_backend(oprt_backend_t backend);
-OPRT_EXPORT oprt_status_t oprt_get_backend(oprt_backend_t *backend);
-
-#ifdef __cplusplus
-}
-#endif
--- a/include/operator_runtime/descriptor.h
+++ b/include/operator_runtime/descriptor.h
@ -1,16 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-
-#include <string>
-
-struct oprt_operator_descriptor {
-    virtual ~oprt_operator_descriptor() = default;
-    virtual const char *op_name() const = 0;
-    size_t workspace_size = 0;
-    oprt_backend_t backend = OPRT_BACKEND_AUTO;
-};
-
-#endif
--- a/include/operator_runtime/detail/cuda_helpers.h
+++ b/include/operator_runtime/detail/cuda_helpers.h
@ -1,29 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#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; \
-        } \
-    } while (0)
-
-namespace oprt {
-
-inline cudaStream_t as_cuda_stream(oprt_stream_t stream) {
-    return reinterpret_cast<cudaStream_t>(stream);
-}
-
-inline int blocks_for(int64_t n, int threads) {
-    int64_t blocks = (n + threads - 1) / threads;
-    return static_cast<int>(blocks > 0 ? blocks : 1);
-}
-
-} // namespace oprt
-
-#endif
--- a/include/operator_runtime/detail/elementwise.h
+++ b/include/operator_runtime/detail/elementwise.h
@ -1,249 +0,0 @@
-#pragma once
-
-#include "operator_runtime/descriptor.h"
-#include "operator_runtime/tensor_view.h"
-#include "operator_runtime/detail/tensor_checks.h"
-
-#ifdef __cplusplus
-
-#include <array>
-#include <cstddef>
-#include <cstdint>
-#include <utility>
-#include <vector>
-
-namespace oprt {
-
-struct ElementwiseInfo {
-    struct InputInfo {
-        std::array<int64_t, OPRT_MAX_DIMS> shape{};
-        std::array<int64_t, OPRT_MAX_DIMS> strides{};
-        bool contiguous = false;
-        bool broadcasted = false;
-    };
-
-    oprt_dtype_t dtype = OPRT_DTYPE_F32;
-    int32_t ndim = 0;
-    int64_t elements = 0;
-    size_t input_count = 0;
-    std::array<int64_t, OPRT_MAX_DIMS> output_shape{};
-    std::array<int64_t, OPRT_MAX_DIMS> output_strides{};
-    bool output_contiguous = false;
-    std::vector<InputInfo> inputs;
-
-    size_t meta_bytes() const {
-        return sizeof(DeviceMeta) + input_count * sizeof(DeviceInputMeta);
-    }
-
-    size_t workspace_bytes() const {
-        return input_count * sizeof(const void *) + meta_bytes();
-    }
-
-    struct DeviceInputMeta {
-        int64_t shape[OPRT_MAX_DIMS];
-        int64_t strides[OPRT_MAX_DIMS];
-        uint8_t contiguous;
-        uint8_t broadcasted;
-    };
-
-    struct DeviceMeta {
-        int32_t ndim;
-        uint32_t input_count;
-        int64_t elements;
-        int64_t output_shape[OPRT_MAX_DIMS];
-        int64_t output_strides[OPRT_MAX_DIMS];
-        uint8_t output_contiguous;
-    };
-};
-
-inline bool elementwise_fast_path(const oprt_tensor_view_t &out,
-                                  const oprt_tensor_view_t &a) {
-    return same_shape(out, a) && is_contiguous(out) && is_contiguous(a);
-}
-
-inline bool elementwise_fast_path(const oprt_tensor_view_t &out,
-                                  const oprt_tensor_view_t &a,
-                                  const oprt_tensor_view_t &b) {
-    return same_shape(out, a) && same_shape(out, b) &&
-           is_contiguous(out) && is_contiguous(a) && is_contiguous(b);
-}
-
-inline bool elementwise_input_broadcastable_to(const oprt_tensor_view_t &out,
-                                               const oprt_tensor_view_t &in) {
-    if (out.ndim < 0 || out.ndim > OPRT_MAX_DIMS || in.ndim < 0 || in.ndim > OPRT_MAX_DIMS) {
-        return false;
-    }
-    if (in.ndim > out.ndim) {
-        return false;
-    }
-    int32_t dim_offset = out.ndim - in.ndim;
-    for (int32_t out_dim = 0; out_dim < out.ndim; ++out_dim) {
-        int32_t in_dim = out_dim - dim_offset;
-        if (in_dim < 0) {
-            continue;
-        }
-        int64_t in_extent = in.shape[in_dim];
-        int64_t out_extent = out.shape[out_dim];
-        if (in_extent != out_extent && in_extent != 1) {
-            return false;
-        }
-    }
-    return true;
-}
-
-inline bool elementwise_has_broadcast(const oprt_tensor_view_t &out,
-                                      const oprt_tensor_view_t &in) {
-    if (has_broadcast_dim(in)) {
-        return true;
-    }
-    if (in.ndim != out.ndim) {
-        return true;
-    }
-    for (int32_t i = 0; i < out.ndim; ++i) {
-        if (in.shape[i] != out.shape[i]) {
-            return true;
-        }
-    }
-    return false;
-}
-
-inline bool elementwise_same_shape_inputs(const oprt_tensor_view_t &out,
-                                          const std::vector<const oprt_tensor_view_t *> &inputs) {
-    for (const auto *input : inputs) {
-        if (input == nullptr || !same_shape(out, *input)) {
-            return false;
-        }
-    }
-    return true;
-}
-
-inline oprt_status_t create_elementwise_info(const oprt_tensor_view_t *out,
-                                             const std::vector<const oprt_tensor_view_t *> &inputs,
-                                             ElementwiseInfo *info) {
-    if (info == nullptr || out == nullptr || inputs.empty()) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-
-    auto status = check_tensor(out);
-    if (status != OPRT_SUCCESS) {
-        return status;
-    }
-    if (has_broadcast_dim(*out)) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-
-    ElementwiseInfo next;
-    next.dtype = out->dtype;
-    next.ndim = out->ndim;
-    next.elements = numel(*out);
-    next.input_count = inputs.size();
-    next.output_contiguous = is_contiguous(*out);
-
-    for (int32_t i = 0; i < out->ndim; ++i) {
-        next.output_shape[i] = out->shape[i];
-        next.output_strides[i] = out->strides[i];
-    }
-
-    next.inputs.reserve(inputs.size());
-    for (const auto *input : inputs) {
-        status = check_tensor(input);
-        if (status != OPRT_SUCCESS) {
-            return status;
-        }
-        if (input->dtype != out->dtype) {
-            return OPRT_ERR_UNSUPPORTED_DTYPE;
-        }
-        if (!elementwise_input_broadcastable_to(*out, *input)) {
-            return OPRT_ERR_INVALID_ARG;
-        }
-
-        ElementwiseInfo::InputInfo input_info;
-        input_info.contiguous = same_shape(*out, *input) && is_contiguous(*input);
-        input_info.broadcasted = elementwise_has_broadcast(*out, *input);
-
-        int32_t dim_offset = out->ndim - input->ndim;
-        for (int32_t out_dim = 0; out_dim < out->ndim; ++out_dim) {
-            int32_t in_dim = out_dim - dim_offset;
-            if (in_dim < 0) {
-                input_info.shape[out_dim] = out->shape[out_dim];
-                input_info.strides[out_dim] = 0;
-                input_info.broadcasted = true;
-                continue;
-            }
-            input_info.shape[out_dim] = out->shape[out_dim];
-            input_info.strides[out_dim] = input->shape[in_dim] == 1 && out->shape[out_dim] != 1
-                                             ? 0
-                                             : input->strides[in_dim];
-        }
-        next.inputs.push_back(input_info);
-    }
-
-    *info = std::move(next);
-    return OPRT_SUCCESS;
-}
-
-struct ElementwiseDescriptorBase : oprt_operator_descriptor {
-    oprt_tensor_view_t out_view{};
-    std::vector<oprt_tensor_view_t> input_views;
-    oprt::ElementwiseInfo info;
-};
-
-inline oprt_status_t init_elementwise_descriptor(ElementwiseDescriptorBase *desc,
-                                                 const oprt_tensor_view_t *out,
-                                                 const std::vector<const oprt_tensor_view_t *> &inputs) {
-    if (desc == nullptr || out == nullptr || inputs.empty()) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-
-    oprt::ElementwiseInfo info;
-    auto status = oprt::create_elementwise_info(out, inputs, &info);
-    if (status != OPRT_SUCCESS) {
-        return status;
-    }
-
-    desc->out_view = *out;
-    desc->input_views.clear();
-    desc->input_views.reserve(inputs.size());
-    for (const auto *input : inputs) {
-        desc->input_views.push_back(*input);
-    }
-    desc->info = std::move(info);
-    desc->workspace_size = desc->info.workspace_bytes();
-    return OPRT_SUCCESS;
-}
-
-inline oprt_status_t get_elementwise_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;
-}
-
-inline oprt_status_t validate_elementwise_execute_args(oprt_operator_descriptor_t desc,
-                                                       size_t workspace_size,
-                                                       void *out,
-                                                       const std::vector<const void *> &inputs) {
-    if (desc == nullptr || out == nullptr || inputs.empty()) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    for (const auto *input : inputs) {
-        if (input == nullptr) {
-            return OPRT_ERR_INVALID_ARG;
-        }
-    }
-    if (workspace_size < desc->workspace_size) {
-        return OPRT_ERR_INSUFFICIENT_WORKSPACE;
-    }
-    return OPRT_SUCCESS;
-}
-
-inline oprt_status_t destroy_elementwise_descriptor(oprt_operator_descriptor_t desc) {
-    delete desc;
-    return OPRT_SUCCESS;
-}
-
-} // namespace oprt
-
-#endif
--- a/include/operator_runtime/detail/operation.h
+++ b/include/operator_runtime/detail/operation.h
@ -1,13 +0,0 @@
-#pragma once
-
-#include <string>
-
-namespace oprt {
-
-struct OperationSpec {
-    std::string name;
-    std::string backend;
-    std::string kind;
-};
-
-} // namespace oprt
--- a/include/operator_runtime/detail/tensor_checks.h
+++ b/include/operator_runtime/detail/tensor_checks.h
@ -1,39 +0,0 @@
-#pragma once
-
-#include "operator_runtime/tensor_view.h"
-
-#ifdef __cplusplus
-
-namespace oprt {
-
-inline oprt_status_t check_tensor(const oprt_tensor_view_t *view) {
-    if (view == nullptr || view->data == nullptr) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    if (view->ndim < 0 || view->ndim > OPRT_MAX_DIMS) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    for (int32_t i = 0; i < view->ndim; ++i) {
-        if (view->shape[i] < 0) {
-            return OPRT_ERR_INVALID_ARG;
-        }
-    }
-    if (view->dtype != OPRT_DTYPE_F16 && view->dtype != OPRT_DTYPE_F32) {
-        return OPRT_ERR_UNSUPPORTED_DTYPE;
-    }
-    return OPRT_SUCCESS;
-}
-
-inline oprt_status_t check_same_dtype(const oprt_tensor_view_t &a,
-                                      const oprt_tensor_view_t &b) {
-    return a.dtype == b.dtype ? OPRT_SUCCESS : OPRT_ERR_UNSUPPORTED_DTYPE;
-}
-
-inline oprt_status_t check_same_shape(const oprt_tensor_view_t &a,
-                                      const oprt_tensor_view_t &b) {
-    return same_shape(a, b) ? OPRT_SUCCESS : OPRT_ERR_INVALID_ARG;
-}
-
-} // namespace oprt
-
-#endif
--- a/include/operator_runtime/operator_runtime.h
+++ b/include/operator_runtime/operator_runtime.h
@ -1,10 +0,0 @@
-#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/relu.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
@ -1,31 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-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(
-    oprt_operator_descriptor_t desc,
-    size_t *size);
-
-OPRT_EXPORT oprt_status_t oprt_execute_copy(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *dst,
-    const void *src,
-    oprt_stream_t stream);
-
-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
@ -1,32 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-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_reduce_sum_workspace_size(
-    oprt_operator_descriptor_t desc,
-    size_t *size);
-
-OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *out,
-    const void *in,
-    oprt_stream_t stream);
-
-OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor(
-    oprt_operator_descriptor_t desc);
-
-#ifdef __cplusplus
-}
-#endif
--- a/include/operator_runtime/ops/relu.h
+++ b/include/operator_runtime/ops/relu.h
@ -1,32 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-OPRT_EXPORT oprt_status_t oprt_create_relu_descriptor(
-    oprt_operator_descriptor_t *desc,
-    const oprt_tensor_view_t *out,
-    const oprt_tensor_view_t *in,
-    float negative_slope);
-
-OPRT_EXPORT oprt_status_t oprt_get_relu_workspace_size(
-    oprt_operator_descriptor_t desc,
-    size_t *size);
-
-OPRT_EXPORT oprt_status_t oprt_execute_relu(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *out,
-    const void *in,
-    oprt_stream_t stream);
-
-OPRT_EXPORT oprt_status_t oprt_destroy_relu_descriptor(
-    oprt_operator_descriptor_t desc);
-
-#ifdef __cplusplus
-}
-#endif
--- a/include/operator_runtime/ops/softmax.h
+++ b/include/operator_runtime/ops/softmax.h
@ -1,32 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-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_softmax_workspace_size(
-    oprt_operator_descriptor_t desc,
-    size_t *size);
-
-OPRT_EXPORT oprt_status_t oprt_execute_softmax(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *out,
-    const void *in,
-    oprt_stream_t stream);
-
-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
@ -1,33 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-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(
-    oprt_operator_descriptor_t desc,
-    size_t *size);
-
-OPRT_EXPORT oprt_status_t oprt_execute_vector_add(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *out,
-    const void *a,
-    const void *b,
-    oprt_stream_t stream);
-
-OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor(
-    oprt_operator_descriptor_t desc);
-
-#ifdef __cplusplus
-}
-#endif
--- a/include/operator_runtime/tensor_view.h
+++ b/include/operator_runtime/tensor_view.h
@ -1,66 +0,0 @@
-#pragma once
-
-#include "operator_runtime/api.h"
-
-#ifdef __cplusplus
-
-#include <algorithm>
-#include <cstdint>
-
-namespace oprt {
-
-inline int64_t numel(const oprt_tensor_view_t &view) {
-    if (view.ndim < 0 || view.ndim > OPRT_MAX_DIMS) {
-        return 0;
-    }
-    int64_t total = 1;
-    for (int32_t i = 0; i < view.ndim; ++i) {
-        total *= view.shape[i];
-    }
-    return total;
-}
-
-inline bool is_contiguous(const oprt_tensor_view_t &view) {
-    if (view.ndim < 0 || view.ndim > OPRT_MAX_DIMS) {
-        return false;
-    }
-    int64_t expected = 1;
-    for (int32_t i = view.ndim - 1; i >= 0; --i) {
-        if (view.shape[i] == 1) {
-            continue;
-        }
-        if (view.strides[i] != expected) {
-            return false;
-        }
-        expected *= view.shape[i];
-    }
-    return true;
-}
-
-inline bool same_shape(const oprt_tensor_view_t &a, const oprt_tensor_view_t &b) {
-    if (a.ndim != b.ndim) {
-        return false;
-    }
-    for (int32_t i = 0; i < a.ndim; ++i) {
-        if (a.shape[i] != b.shape[i]) {
-            return false;
-        }
-    }
-    return true;
-}
-
-inline bool has_broadcast_dim(const oprt_tensor_view_t &view) {
-    if (view.ndim < 0 || view.ndim > OPRT_MAX_DIMS) {
-        return false;
-    }
-    for (int32_t i = 0; i < view.ndim; ++i) {
-        if (view.shape[i] > 1 && view.strides[i] == 0) {
-            return true;
-        }
-    }
-    return false;
-}
-
-} // namespace oprt
-
-#endif
--- a/ops/CMakeLists.txt
+++ b/ops/CMakeLists.txt
@ -1,59 +0,0 @@
-# Convention: operator CUDA sources auto-discovered by directory layout:
-#   ops/<operator>/nvidia/*.cu  ->  compiled when CAMP_ENABLE_NVIDIA is ON
-#   ops/elementwise/<operator>/nvidia/*.cu  ->  compiled when CAMP_ENABLE_NVIDIA is ON
-#   ops/<operator>/metax/*.maca ->  compiled when CAMP_ENABLE_METAX is ON
-#   ops/elementwise/<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
-  ${CMAKE_CURRENT_SOURCE_DIR}/common/status.cc
-)
-
-set(CAMP_NVIDIA_SOURCES "")
-if(CAMP_ENABLE_NVIDIA)
-  file(GLOB_RECURSE CAMP_NVIDIA_SOURCES
-    ${CMAKE_CURRENT_SOURCE_DIR}/*/nvidia/*.cu
-    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise/*/nvidia/*.cu
-  )
-endif()
-
-set(CAMP_METAX_SOURCES "")
-if(CAMP_ENABLE_METAX)
-  file(GLOB_RECURSE CAMP_METAX_SOURCES
-    ${CMAKE_CURRENT_SOURCE_DIR}/*/metax/*.maca
-    ${CMAKE_CURRENT_SOURCE_DIR}/elementwise/*/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}/../include"
-  PRIVATE
-    "${CMAKE_CURRENT_SOURCE_DIR}/.."
-)
-
-if(CAMP_USE_CUTE)
-  target_include_directories(camp_ops PRIVATE ${CAMP_CUTE_TARGET_INCLUDE_DIRS})
-  target_compile_definitions(camp_ops PRIVATE CAMP_ENABLE_CUTE=1)
-  target_compile_options(camp_ops PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>)
-endif()
-
-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>
-)
-
-if(CAMP_ENABLE_NVIDIA)
-  set_target_properties(camp_ops PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
-endif()
--- a/ops/common/elementwise/metax/elementwise_metax.h
+++ b/ops/common/elementwise/metax/elementwise_metax.h
@ -1,184 +0,0 @@
-#pragma once
-
-#include "operator_runtime/detail/cuda_helpers.h"
-#include "operator_runtime/detail/elementwise.h"
-
-#include <cuda_fp16.h>
-#include <cuda_runtime.h>
-
-#include <array>
-#include <cstdint>
-#include <utility>
-
-namespace oprt::elementwise::metax {
-
-template <typename T>
-__device__ inline T cast_scalar(float value) {
-    return static_cast<T>(value);
-}
-
-template <>
-__device__ inline half cast_scalar<half>(float value) {
-    return __float2half(value);
-}
-
-__device__ inline int64_t index_to_offset(int64_t linear,
-                                          int32_t ndim,
-                                          const int64_t *shape,
-                                          const int64_t *strides) {
-    int64_t offset = 0;
-    for (int32_t dim = ndim - 1; dim >= 0; --dim) {
-        int64_t coord = linear % shape[dim];
-        linear /= shape[dim];
-        offset += coord * strides[dim];
-    }
-    return offset;
-}
-
-template <size_t Input>
-__device__ inline int64_t input_offset(
-    int64_t linear,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta) {
-    const auto &m = input_meta[Input];
-    return m.contiguous != 0
-             ? linear
-             : index_to_offset(linear, meta->ndim, m.shape, m.strides);
-}
-
-template <typename T, size_t Input>
-__device__ inline T load_input(
-    int64_t linear,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta) {
-    const T *input = static_cast<const T *>(inputs[Input]);
-    return input[input_offset<Input>(linear, meta, input_meta)];
-}
-
-template <typename T, typename Op, size_t... Is, typename... Args>
-__device__ inline T apply_at_index_impl(
-    int64_t linear,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta,
-    Op op,
-    std::index_sequence<Is...>,
-    Args... args) {
-    static_assert(sizeof...(Is) > 0, "elementwise launch requires at least one input");
-    return op(load_input<T, Is>(linear, inputs, meta, input_meta)..., args...);
-}
-
-template <typename T, size_t N, typename Op, typename... Args>
-__device__ inline T apply_at_index(
-    int64_t linear,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta,
-    Op op,
-    Args... args) {
-    return apply_at_index_impl<T>(
-        linear,
-        inputs,
-        meta,
-        input_meta,
-        op,
-        std::make_index_sequence<N>{},
-        args...);
-}
-
-template <typename T, size_t N, typename Op, typename... Args>
-__global__ void elementwise_kernel(
-    T *__restrict__ out,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta,
-    Op op,
-    Args... args) {
-    int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
-    int64_t stride = static_cast<int64_t>(blockDim.x) * gridDim.x;
-    const bool output_contiguous = meta->output_contiguous != 0;
-
-    for (int64_t i = idx; i < meta->elements; i += stride) {
-        int64_t out_offset = output_contiguous
-                               ? i
-                               : index_to_offset(i, meta->ndim, meta->output_shape, meta->output_strides);
-        out[out_offset] = apply_at_index<T, N>(i, inputs, meta, input_meta, op, args...);
-    }
-}
-
-template <typename T, size_t N, typename Op, typename... Args>
-oprt_status_t launch(
-    const oprt::ElementwiseInfo &info,
-    void *workspace,
-    void *out,
-    const std::array<const void *, N> &inputs,
-    oprt_stream_t stream,
-    Op op,
-    Args... args) {
-    if (info.input_count != N || (info.workspace_bytes() != 0 && workspace == nullptr)) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    if (info.elements == 0) {
-        return OPRT_SUCCESS;
-    }
-
-    auto *workspace_bytes = static_cast<uint8_t *>(workspace);
-    auto **device_inputs = reinterpret_cast<const void **>(workspace_bytes);
-    auto *device_meta = reinterpret_cast<oprt::ElementwiseInfo::DeviceMeta *>(
-        workspace_bytes + N * sizeof(const void *));
-    auto *device_input_meta = reinterpret_cast<oprt::ElementwiseInfo::DeviceInputMeta *>(device_meta + 1);
-
-    oprt::ElementwiseInfo::DeviceMeta host_meta{};
-    host_meta.ndim = info.ndim;
-    host_meta.input_count = static_cast<uint32_t>(N);
-    host_meta.elements = info.elements;
-    host_meta.output_contiguous = info.output_contiguous ? 1 : 0;
-    for (int32_t i = 0; i < info.ndim; ++i) {
-        host_meta.output_shape[i] = info.output_shape[i];
-        host_meta.output_strides[i] = info.output_strides[i];
-    }
-
-    std::array<oprt::ElementwiseInfo::DeviceInputMeta, N> host_input_meta{};
-    for (size_t input = 0; input < N; ++input) {
-        const auto &src = info.inputs[input];
-        auto &dst = host_input_meta[input];
-        dst.contiguous = src.contiguous ? 1 : 0;
-        dst.broadcasted = src.broadcasted ? 1 : 0;
-        for (int32_t dim = 0; dim < info.ndim; ++dim) {
-            dst.shape[dim] = src.shape[dim];
-            dst.strides[dim] = src.strides[dim];
-        }
-    }
-
-    cudaStream_t cuda_stream = oprt::as_cuda_stream(stream);
-    OPRT_CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(device_inputs,
-                                              inputs.data(),
-                                              N * sizeof(const void *),
-                                              cudaMemcpyHostToDevice,
-                                              cuda_stream));
-    OPRT_CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(device_meta,
-                                              &host_meta,
-                                              sizeof(host_meta),
-                                              cudaMemcpyHostToDevice,
-                                              cuda_stream));
-    OPRT_CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(device_input_meta,
-                                              host_input_meta.data(),
-                                              N * sizeof(oprt::ElementwiseInfo::DeviceInputMeta),
-                                              cudaMemcpyHostToDevice,
-                                              cuda_stream));
-
-    constexpr int threads = 256;
-    int blocks = oprt::blocks_for(info.elements, threads);
-    elementwise_kernel<T, N, Op, Args...><<<blocks, threads, 0, cuda_stream>>>(
-        static_cast<T *>(out),
-        device_inputs,
-        device_meta,
-        device_input_meta,
-        op,
-        args...);
-    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
-    return OPRT_SUCCESS;
-}
-
-} // namespace oprt::elementwise::metax
--- a/ops/common/elementwise/nvidia/elementwise_nvidia.cuh
+++ b/ops/common/elementwise/nvidia/elementwise_nvidia.cuh
@ -1,184 +0,0 @@
-#pragma once
-
-#include "operator_runtime/detail/cuda_helpers.h"
-#include "operator_runtime/detail/elementwise.h"
-
-#include <cuda_fp16.h>
-#include <cuda_runtime.h>
-
-#include <array>
-#include <cstdint>
-#include <utility>
-
-namespace oprt::elementwise::nvidia {
-
-template <typename T>
-__device__ inline T cast_scalar(float value) {
-    return static_cast<T>(value);
-}
-
-template <>
-__device__ inline half cast_scalar<half>(float value) {
-    return __float2half(value);
-}
-
-__device__ inline int64_t index_to_offset(int64_t linear,
-                                          int32_t ndim,
-                                          const int64_t *shape,
-                                          const int64_t *strides) {
-    int64_t offset = 0;
-    for (int32_t dim = ndim - 1; dim >= 0; --dim) {
-        int64_t coord = linear % shape[dim];
-        linear /= shape[dim];
-        offset += coord * strides[dim];
-    }
-    return offset;
-}
-
-template <size_t Input>
-__device__ inline int64_t input_offset(
-    int64_t linear,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta) {
-    const auto &m = input_meta[Input];
-    return m.contiguous != 0
-             ? linear
-             : index_to_offset(linear, meta->ndim, m.shape, m.strides);
-}
-
-template <typename T, size_t Input>
-__device__ inline T load_input(
-    int64_t linear,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta) {
-    const T *input = static_cast<const T *>(inputs[Input]);
-    return input[input_offset<Input>(linear, meta, input_meta)];
-}
-
-template <typename T, typename Op, size_t... Is, typename... Args>
-__device__ inline T apply_at_index_impl(
-    int64_t linear,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta,
-    Op op,
-    std::index_sequence<Is...>,
-    Args... args) {
-    static_assert(sizeof...(Is) > 0, "elementwise launch requires at least one input");
-    return op(load_input<T, Is>(linear, inputs, meta, input_meta)..., args...);
-}
-
-template <typename T, size_t N, typename Op, typename... Args>
-__device__ inline T apply_at_index(
-    int64_t linear,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta,
-    Op op,
-    Args... args) {
-    return apply_at_index_impl<T>(
-        linear,
-        inputs,
-        meta,
-        input_meta,
-        op,
-        std::make_index_sequence<N>{},
-        args...);
-}
-
-template <typename T, size_t N, typename Op, typename... Args>
-__global__ void elementwise_kernel(
-    T *__restrict__ out,
-    const void *const *__restrict__ inputs,
-    const oprt::ElementwiseInfo::DeviceMeta *__restrict__ meta,
-    const oprt::ElementwiseInfo::DeviceInputMeta *__restrict__ input_meta,
-    Op op,
-    Args... args) {
-    int64_t idx = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
-    int64_t stride = static_cast<int64_t>(blockDim.x) * gridDim.x;
-    const bool output_contiguous = meta->output_contiguous != 0;
-
-    for (int64_t i = idx; i < meta->elements; i += stride) {
-        int64_t out_offset = output_contiguous
-                               ? i
-                               : index_to_offset(i, meta->ndim, meta->output_shape, meta->output_strides);
-        out[out_offset] = apply_at_index<T, N>(i, inputs, meta, input_meta, op, args...);
-    }
-}
-
-template <typename T, size_t N, typename Op, typename... Args>
-oprt_status_t launch(
-    const oprt::ElementwiseInfo &info,
-    void *workspace,
-    void *out,
-    const std::array<const void *, N> &inputs,
-    oprt_stream_t stream,
-    Op op,
-    Args... args) {
-    if (info.input_count != N || (info.workspace_bytes() != 0 && workspace == nullptr)) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    if (info.elements == 0) {
-        return OPRT_SUCCESS;
-    }
-
-    auto *workspace_bytes = static_cast<uint8_t *>(workspace);
-    auto **device_inputs = reinterpret_cast<const void **>(workspace_bytes);
-    auto *device_meta = reinterpret_cast<oprt::ElementwiseInfo::DeviceMeta *>(
-        workspace_bytes + N * sizeof(const void *));
-    auto *device_input_meta = reinterpret_cast<oprt::ElementwiseInfo::DeviceInputMeta *>(device_meta + 1);
-
-    oprt::ElementwiseInfo::DeviceMeta host_meta{};
-    host_meta.ndim = info.ndim;
-    host_meta.input_count = static_cast<uint32_t>(N);
-    host_meta.elements = info.elements;
-    host_meta.output_contiguous = info.output_contiguous ? 1 : 0;
-    for (int32_t i = 0; i < info.ndim; ++i) {
-        host_meta.output_shape[i] = info.output_shape[i];
-        host_meta.output_strides[i] = info.output_strides[i];
-    }
-
-    std::array<oprt::ElementwiseInfo::DeviceInputMeta, N> host_input_meta{};
-    for (size_t input = 0; input < N; ++input) {
-        const auto &src = info.inputs[input];
-        auto &dst = host_input_meta[input];
-        dst.contiguous = src.contiguous ? 1 : 0;
-        dst.broadcasted = src.broadcasted ? 1 : 0;
-        for (int32_t dim = 0; dim < info.ndim; ++dim) {
-            dst.shape[dim] = src.shape[dim];
-            dst.strides[dim] = src.strides[dim];
-        }
-    }
-
-    cudaStream_t cuda_stream = oprt::as_cuda_stream(stream);
-    OPRT_CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(device_inputs,
-                                              inputs.data(),
-                                              N * sizeof(const void *),
-                                              cudaMemcpyHostToDevice,
-                                              cuda_stream));
-    OPRT_CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(device_meta,
-                                              &host_meta,
-                                              sizeof(host_meta),
-                                              cudaMemcpyHostToDevice,
-                                              cuda_stream));
-    OPRT_CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(device_input_meta,
-                                              host_input_meta.data(),
-                                              N * sizeof(oprt::ElementwiseInfo::DeviceInputMeta),
-                                              cudaMemcpyHostToDevice,
-                                              cuda_stream));
-
-    constexpr int threads = 256;
-    int blocks = oprt::blocks_for(info.elements, threads);
-    elementwise_kernel<T, N, Op, Args...><<<blocks, threads, 0, cuda_stream>>>(
-        static_cast<T *>(out),
-        device_inputs,
-        device_meta,
-        device_input_meta,
-        op,
-        args...);
-    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
-    return OPRT_SUCCESS;
-}
-
-} // namespace oprt::elementwise::nvidia
--- a/ops/common/status.cc
+++ b/ops/common/status.cc
@ -1,82 +0,0 @@
-#include "operator_runtime/api.h"
-
-namespace {
-#ifdef CAMP_ENABLE_NVIDIA
-oprt_backend_t g_current_backend = OPRT_BACKEND_NVIDIA;
-#elif defined(CAMP_ENABLE_METAX)
-oprt_backend_t g_current_backend = OPRT_BACKEND_METAX;
-#else
-oprt_backend_t g_current_backend = OPRT_BACKEND_AUTO;
-#endif
-}
-
-extern "C" OPRT_EXPORT const char *oprt_status_string(oprt_status_t status) {
-    switch (status) {
-    case OPRT_SUCCESS:
-        return "success";
-    case OPRT_ERR_INVALID_ARG:
-        return "invalid argument";
-    case OPRT_ERR_UNSUPPORTED_DTYPE:
-        return "unsupported dtype";
-    case OPRT_ERR_RUNTIME:
-        return "runtime error";
-    case OPRT_ERR_INSUFFICIENT_WORKSPACE:
-        return "insufficient workspace";
-    case OPRT_ERR_NOT_SUPPORTED:
-        return "not supported";
-    default:
-        return "unknown status";
-    }
-}
-
-extern "C" OPRT_EXPORT const char *oprt_backend_string(oprt_backend_t backend) {
-    switch (backend) {
-    case OPRT_BACKEND_AUTO:
-        return "auto";
-    case OPRT_BACKEND_NVIDIA:
-        return "nvidia";
-    case OPRT_BACKEND_METAX:
-        return "metax";
-    default:
-        return "unknown";
-    }
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_set_backend(oprt_backend_t backend) {
-    switch (backend) {
-    case OPRT_BACKEND_AUTO:
-#ifdef CAMP_ENABLE_NVIDIA
-        g_current_backend = OPRT_BACKEND_NVIDIA;
-        return OPRT_SUCCESS;
-#elif defined(CAMP_ENABLE_METAX)
-        g_current_backend = OPRT_BACKEND_METAX;
-        return OPRT_SUCCESS;
-#else
-        return OPRT_ERR_NOT_SUPPORTED;
-#endif
-    case OPRT_BACKEND_NVIDIA:
-#ifdef CAMP_ENABLE_NVIDIA
-        g_current_backend = backend;
-        return OPRT_SUCCESS;
-#else
-        return OPRT_ERR_NOT_SUPPORTED;
-#endif
-    case OPRT_BACKEND_METAX:
-#ifdef CAMP_ENABLE_METAX
-        g_current_backend = backend;
-        return OPRT_SUCCESS;
-#else
-        return OPRT_ERR_NOT_SUPPORTED;
-#endif
-    default:
-        return OPRT_ERR_INVALID_ARG;
-    }
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_get_backend(oprt_backend_t *backend) {
-    if (backend == nullptr) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    *backend = g_current_backend;
-    return OPRT_SUCCESS;
-}
--- a/ops/common/tilelang/eager_copy.py
+++ b/ops/common/tilelang/eager_copy.py
@ -1,86 +0,0 @@
-from dataclasses import dataclass
-from functools import lru_cache
-
-import tilelang
-import tilelang.language as T
-import torch
-
-
-def _tl_dtype(dtype: torch.dtype):
-    if dtype is torch.float16:
-        return T.float16
-    if dtype is torch.float32:
-        return T.float32
-    raise TypeError(f"unsupported TileLang dtype: {dtype}")
-
-
-@tilelang.jit
-def _copy_eager_kernel(src, BLOCK_N: int, dtype):
-    N = T.const("N")
-    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],
-        )
-
-    return out
-
-
-def _block_n(n: int) -> int:
-    return 1024 if n % 1024 == 0 else n
-
-
-@lru_cache(maxsize=32)
-def _compiled_copy_eager(n: int, block_n: int, dtype: torch.dtype):
-    return _copy_eager_kernel.compile(N=n, BLOCK_N=block_n, dtype=_tl_dtype(dtype))
-
-
-@dataclass
-class EagerCopyPrepared:
-    out: torch.Tensor
-    src: torch.Tensor
-    kernel: object
-
-    def run(self) -> None:
-        self.out.copy_(self.kernel(self.src))
-
-    def destroy(self) -> None:
-        pass
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-
-def prepare_copy_eager(out: torch.Tensor, src: torch.Tensor) -> EagerCopyPrepared:
-    if out.shape != src.shape:
-        raise ValueError("eager copy expects matching shapes")
-    if out.dtype != src.dtype:
-        raise TypeError("eager copy expects matching dtypes")
-    if not out.is_cuda or not src.is_cuda:
-        raise ValueError("eager copy expects CUDA tensors")
-    if not out.is_contiguous() or not src.is_contiguous():
-        raise ValueError("eager copy v1 supports contiguous tensors only")
-
-    n = src.numel()
-    block_n = _block_n(n)
-    if n % block_n != 0:
-        raise ValueError("eager copy v1 requires N % BLOCK_N == 0")
-    kernel = _compiled_copy_eager(n, block_n, src.dtype)
-    return EagerCopyPrepared(out, src, kernel)
-
-
-def copy_eager_(out: torch.Tensor, src: torch.Tensor) -> torch.Tensor:
-    with prepare_copy_eager(out, src) as prepared:
-        prepared.run()
-    return out
-
-
-def copy_eager(src: torch.Tensor) -> torch.Tensor:
-    out = torch.empty_like(src)
-    return copy_eager_(out, src)
--- a/ops/common/tilelang/lazy_out_idx_copy.py
+++ b/ops/common/tilelang/lazy_out_idx_copy.py
@ -1,91 +0,0 @@
-from dataclasses import dataclass
-from functools import lru_cache
-
-import tilelang
-import tilelang.language as T
-import torch
-
-
-def _tl_dtype_str(dtype: torch.dtype) -> str:
-    if dtype is torch.float16:
-        return "float16"
-    if dtype is torch.float32:
-        return "float32"
-    raise TypeError(f"unsupported TileLang dtype: {dtype}")
-
-
-@tilelang.jit(out_idx=[1])
-def _copy_lazy_out_idx_kernel(n: int, block_n: int, dtype: str):
-    @T.prim_func
-    def main(src: T.Tensor((n,), dtype), out: T.Tensor((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],
-            )
-
-    return main
-
-
-def _block_n(n: int) -> int:
-    return 1024 if n % 1024 == 0 else n
-
-
-@lru_cache(maxsize=32)
-def _compiled_copy_lazy_out_idx(n: int, block_n: int, dtype: torch.dtype):
-    return _copy_lazy_out_idx_kernel(n, block_n, _tl_dtype_str(dtype))
-
-
-@dataclass
-class LazyOutIdxCopyPrepared:
-    out: torch.Tensor
-    src: torch.Tensor
-    kernel: object
-
-    def run(self) -> None:
-        # out_idx makes TileLang allocate and return the output tensor. The
-        # training runtime keeps an out-variant API, so this adapter copies the
-        # result into the caller-owned output buffer.
-        self.out.copy_(self.kernel(self.src))
-
-    def destroy(self) -> None:
-        pass
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-
-def prepare_copy_lazy_out_idx(out: torch.Tensor, src: torch.Tensor) -> LazyOutIdxCopyPrepared:
-    if out.shape != src.shape:
-        raise ValueError("lazy out_idx copy expects matching shapes")
-    if out.dtype != src.dtype:
-        raise TypeError("lazy out_idx copy expects matching dtypes")
-    if not out.is_cuda or not src.is_cuda:
-        raise ValueError("lazy out_idx copy expects CUDA tensors")
-    if not out.is_contiguous() or not src.is_contiguous():
-        raise ValueError("lazy out_idx copy v1 supports contiguous tensors only")
-
-    n = src.numel()
-    block_n = _block_n(n)
-    if n % block_n != 0:
-        raise ValueError("lazy out_idx copy v1 requires N % BLOCK_N == 0")
-    kernel = _compiled_copy_lazy_out_idx(n, block_n, src.dtype)
-    return LazyOutIdxCopyPrepared(out, src, kernel)
-
-
-def copy_lazy_out_idx_(out: torch.Tensor, src: torch.Tensor) -> torch.Tensor:
-    with prepare_copy_lazy_out_idx(out, src) as prepared:
-        prepared.run()
-    return out
-
-
-def copy_lazy_out_idx(src: torch.Tensor) -> torch.Tensor:
-    n = src.numel()
-    block_n = _block_n(n)
-    if n % block_n != 0:
-        raise ValueError("lazy out_idx copy v1 requires N % BLOCK_N == 0")
-    kernel = _compiled_copy_lazy_out_idx(n, block_n, src.dtype)
-    return kernel(src)
--- a/ops/copy/metax/README.md
+++ b/ops/copy/metax/README.md
@ -1,4 +0,0 @@
-# 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.maca
+++ b/ops/copy/metax/copy_metax.maca
@ -1,119 +0,0 @@
-#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,119 +0,0 @@
-#include "operator_runtime/ops/copy.h"
-
-#include "operator_runtime/descriptor.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"
-#ifdef CAMP_ENABLE_CUTE
-#include "ops/copy/nvidia/cute_copy.cuh"
-#endif
-
-#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>
-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);
-#ifdef CAMP_ENABLE_CUTE
-    oprt::copy::nvidia::launch_copy_contiguous_cute_l40<T>(
-        static_cast<T *>(dst), static_cast<const T *>(src), desc->elements, s);
-#else
-    oprt::copy::nvidia::copy_contiguous_kernel<T><<<blocks, threads, 0, s>>>(
-        static_cast<T *>(dst), static_cast<const T *>(src), desc->elements);
-#endif
-    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/cute_copy.cuh
+++ b/ops/copy/nvidia/cute_copy.cuh
@ -1,79 +0,0 @@
-#pragma once
-
-#include <cute/tensor.hpp>
-#include <cuda_runtime.h>
-
-#include <cstdint>
-
-namespace oprt::copy::nvidia {
-
-template <typename T, int Threads, int ElementsPerAccess>
-__global__ void copy_contiguous_cute_l40_kernel(T *dst, const T *src, int64_t n) {
-    using namespace cute;
-
-    constexpr int elements_per_block = Threads * ElementsPerAccess;
-
-    Tensor src_tensor = make_tensor(make_gmem_ptr(src), make_layout(make_shape(n)));
-    Tensor dst_tensor = make_tensor(make_gmem_ptr(dst), make_layout(make_shape(n)));
-
-    auto block_shape = make_shape(Int<elements_per_block>{});
-    auto block_coord = make_coord(blockIdx.x);
-
-    Tensor coords = make_identity_tensor(shape(src_tensor));
-    Tensor predicates = cute::lazy::transform(coords, [&](auto coord) {
-        return elem_less(coord, shape(src_tensor));
-    });
-
-    Tensor src_tile = local_tile(src_tensor, block_shape, block_coord);
-    Tensor dst_tile = local_tile(dst_tensor, block_shape, block_coord);
-    Tensor pred_tile = local_tile(predicates, block_shape, block_coord);
-
-    Layout thread_layout = make_layout(make_shape(Int<Threads>{}));
-    Layout value_layout = make_layout(make_shape(Int<ElementsPerAccess>{}));
-
-    using AccessType = uint_byte_t<sizeof(T) * ElementsPerAccess>;
-    using CopyOp = UniversalCopy<AccessType>;
-    using Atom = Copy_Atom<CopyOp, T>;
-
-    TiledCopy tiled_copy = make_tiled_copy(Atom{}, thread_layout, value_layout);
-    ThrCopy thread_copy = tiled_copy.get_thread_slice(threadIdx.x);
-
-    Tensor thread_src = thread_copy.partition_S(src_tile);
-    Tensor thread_dst = thread_copy.partition_D(dst_tile);
-    Tensor thread_pred = thread_copy.partition_S(pred_tile);
-    Tensor fragment = make_fragment_like(thread_src);
-
-    copy_if(tiled_copy, thread_pred, thread_src, fragment);
-    copy_if(tiled_copy, thread_pred, fragment, thread_dst);
-}
-
-template <typename T>
-inline void launch_copy_contiguous_cute_l40(T *dst, const T *src, int64_t n, cudaStream_t stream) {
-    constexpr int threads = 256;
-    constexpr int bytes_per_access = 16;
-    constexpr int elements_per_access = bytes_per_access / static_cast<int>(sizeof(T));
-    static_assert(elements_per_access > 0, "copy element type is larger than the vector access");
-
-    if (n <= 0) {
-        return;
-    }
-
-    auto src_addr = reinterpret_cast<std::uintptr_t>(src);
-    auto dst_addr = reinterpret_cast<std::uintptr_t>(dst);
-    bool aligned = (src_addr % bytes_per_access == 0) && (dst_addr % bytes_per_access == 0);
-
-    if (aligned) {
-        constexpr int elements_per_block = threads * elements_per_access;
-        int64_t blocks64 = (n + elements_per_block - 1) / elements_per_block;
-        int blocks = static_cast<int>(blocks64 > 0 ? blocks64 : 1);
-        copy_contiguous_cute_l40_kernel<T, threads, elements_per_access><<<blocks, threads, 0, stream>>>(
-            dst, src, n);
-    } else {
-        int64_t blocks64 = (n + threads - 1) / threads;
-        int blocks = static_cast<int>(blocks64 > 0 ? blocks64 : 1);
-        copy_contiguous_cute_l40_kernel<T, threads, 1><<<blocks, threads, 0, stream>>>(
-            dst, src, n);
-    }
-}
-
-} // namespace oprt::copy::nvidia
--- a/ops/copy/nvidia/kernel.cuh
+++ b/ops/copy/nvidia/kernel.cuh
@ -1,19 +0,0 @@
-#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) {
-    // 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
@ -1,73 +0,0 @@
-from __future__ import annotations
-
-from functools import lru_cache
-from dataclasses import dataclass
-
-import tilelang.language as T
-import torch
-
-from ops.copy.tilelang.kernel import copy_kernel
-
-
-def _tl_dtype(dtype: torch.dtype):
-    if dtype is torch.float16:
-        return T.float16
-    if dtype is torch.float32:
-        return T.float32
-    raise TypeError(f"unsupported TileLang dtype: {dtype}")
-
-def _block_n(n: int) -> int:
-    return 1024 if n % 1024 == 0 else n
-
-
-@lru_cache(maxsize=32)
-def _compiled_copy(n: int, block_n: int, dtype: torch.dtype):
-    return copy_kernel.compile(N=n, BLOCK_N=block_n, dtype=_tl_dtype(dtype))
-
-
-@dataclass
-class TileLangCopyPrepared:
-    out: torch.Tensor
-    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.run_inputs())
-
-    def destroy(self) -> None:
-        pass
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-
-def prepare_copy_tl(out: torch.Tensor, src: torch.Tensor) -> TileLangCopyPrepared:
-    n = src.numel()
-    block_n = _block_n(n)
-    if n % block_n != 0:
-        raise ValueError("TileLang copy v1 requires N % BLOCK_N == 0")
-    kernel = _compiled_copy(n, block_n, src.dtype)
-    return TileLangCopyPrepared(out, src, kernel)
-
-
-def copy_tl_(out: torch.Tensor, src: torch.Tensor) -> torch.Tensor:
-    with prepare_copy_tl(out, src) as prepared:
-        prepared.run()
-    return out
-
-
-def copy_tl(src: torch.Tensor) -> torch.Tensor:
-    out = torch.empty_like(src)
-    return copy_tl_(out, src)
--- a/ops/copy/tilelang/kernel.py
+++ b/ops/copy/tilelang/kernel.py
@ -1,19 +0,0 @@
-from __future__ import annotations
-
-import tilelang
-import tilelang.language as T
-
-
-@tilelang.jit
-def copy_kernel(src, BLOCK_N: int, dtype):
-    N = T.const("N")
-    src: T.Tensor((N,), dtype)
-    out = T.empty((N,), dtype)
-
-    # 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/elementwise/relu/metax/relu_metax.maca
+++ b/ops/elementwise/relu/metax/relu_metax.maca
@ -1,84 +0,0 @@
-#include "operator_runtime/ops/relu.h"
-
-#include "ops/common/elementwise/metax/elementwise_metax.h"
-#include "ops/elementwise/relu/relu_common.h"
-
-#include <array>
-
-namespace {
-
-struct ReluOp {
-    template <typename T>
-    __device__ T operator()(T value, float negative_slope) const {
-        T zero = oprt::elementwise::metax::cast_scalar<T>(0.0f);
-        T slope = oprt::elementwise::metax::cast_scalar<T>(negative_slope);
-        return value > zero ? value : value * slope;
-    }
-};
-
-template <>
-__device__ half ReluOp::operator()<half>(half value, float negative_slope) const {
-    half zero = __float2half(0.0f);
-    half slope = __float2half(negative_slope);
-    return __hgt(value, zero) ? value : __hmul(value, slope);
-}
-
-template <typename T>
-oprt_status_t launch_relu(const oprt::elementwise::ReluDescriptor *desc,
-                          void *workspace,
-                          void *out,
-                          const void *in,
-                          oprt_stream_t stream) {
-    return oprt::elementwise::metax::launch<T, 1>(
-        desc->info,
-        workspace,
-        out,
-        std::array<const void *, 1>{in},
-        stream,
-        ReluOp{},
-        desc->negative_slope);
-}
-
-} // namespace
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_create_relu_descriptor(
-    oprt_operator_descriptor_t *desc,
-    const oprt_tensor_view_t *out,
-    const oprt_tensor_view_t *in,
-    float negative_slope) {
-    return oprt::elementwise::create_relu_descriptor_common(desc, out, in, negative_slope);
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_get_relu_workspace_size(
-    oprt_operator_descriptor_t desc,
-    size_t *size) {
-    return oprt::elementwise::get_relu_workspace_size_common(desc, size);
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_execute_relu(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *out,
-    const void *in,
-    oprt_stream_t stream) {
-    auto status = oprt::elementwise::validate_relu_execute_args(desc, workspace_size, out, in);
-    if (status != OPRT_SUCCESS) {
-        return status;
-    }
-
-    auto *typed = static_cast<const oprt::elementwise::ReluDescriptor *>(desc);
-    switch (typed->out_view.dtype) {
-    case OPRT_DTYPE_F16:
-        return launch_relu<half>(typed, workspace, out, in, stream);
-    case OPRT_DTYPE_F32:
-        return launch_relu<float>(typed, workspace, out, in, stream);
-    default:
-        return OPRT_ERR_UNSUPPORTED_DTYPE;
-    }
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_relu_descriptor(
-    oprt_operator_descriptor_t desc) {
-    return oprt::elementwise::destroy_relu_descriptor_common(desc);
-}
--- a/ops/elementwise/relu/nvidia/relu_cuda.cu
+++ b/ops/elementwise/relu/nvidia/relu_cuda.cu
@ -1,84 +0,0 @@
-#include "operator_runtime/ops/relu.h"
-
-#include "ops/common/elementwise/nvidia/elementwise_nvidia.cuh"
-#include "ops/elementwise/relu/relu_common.h"
-
-#include <array>
-
-namespace {
-
-struct ReluOp {
-    template <typename T>
-    __device__ T operator()(T value, float negative_slope) const {
-        T zero = oprt::elementwise::nvidia::cast_scalar<T>(0.0f);
-        T slope = oprt::elementwise::nvidia::cast_scalar<T>(negative_slope);
-        return value > zero ? value : value * slope;
-    }
-};
-
-template <>
-__device__ half ReluOp::operator()<half>(half value, float negative_slope) const {
-    half zero = __float2half(0.0f);
-    half slope = __float2half(negative_slope);
-    return __hgt(value, zero) ? value : __hmul(value, slope);
-}
-
-template <typename T>
-oprt_status_t launch_relu(const oprt::elementwise::ReluDescriptor *desc,
-                          void *workspace,
-                          void *out,
-                          const void *in,
-                          oprt_stream_t stream) {
-    return oprt::elementwise::nvidia::launch<T, 1>(
-        desc->info,
-        workspace,
-        out,
-        std::array<const void *, 1>{in},
-        stream,
-        ReluOp{},
-        desc->negative_slope);
-}
-
-} // namespace
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_create_relu_descriptor(
-    oprt_operator_descriptor_t *desc,
-    const oprt_tensor_view_t *out,
-    const oprt_tensor_view_t *in,
-    float negative_slope) {
-    return oprt::elementwise::create_relu_descriptor_common(desc, out, in, negative_slope);
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_get_relu_workspace_size(
-    oprt_operator_descriptor_t desc,
-    size_t *size) {
-    return oprt::elementwise::get_relu_workspace_size_common(desc, size);
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_execute_relu(
-    oprt_operator_descriptor_t desc,
-    void *workspace,
-    size_t workspace_size,
-    void *out,
-    const void *in,
-    oprt_stream_t stream) {
-    auto status = oprt::elementwise::validate_relu_execute_args(desc, workspace_size, out, in);
-    if (status != OPRT_SUCCESS) {
-        return status;
-    }
-
-    auto *typed = static_cast<const oprt::elementwise::ReluDescriptor *>(desc);
-    switch (typed->out_view.dtype) {
-    case OPRT_DTYPE_F16:
-        return launch_relu<half>(typed, workspace, out, in, stream);
-    case OPRT_DTYPE_F32:
-        return launch_relu<float>(typed, workspace, out, in, stream);
-    default:
-        return OPRT_ERR_UNSUPPORTED_DTYPE;
-    }
-}
-
-extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_relu_descriptor(
-    oprt_operator_descriptor_t desc) {
-    return oprt::elementwise::destroy_relu_descriptor_common(desc);
-}
--- a/ops/elementwise/relu/relu_common.h
+++ b/ops/elementwise/relu/relu_common.h
@ -1,60 +0,0 @@
-#pragma once
-
-#include "operator_runtime/descriptor.h"
-#include "operator_runtime/detail/elementwise.h"
-
-namespace oprt::elementwise {
-
-struct ReluDescriptor final : oprt::ElementwiseDescriptorBase {
-    float negative_slope = 0.0f;
-
-    const char *op_name() const override {
-        return "relu";
-    }
-};
-
-inline oprt_status_t create_relu_descriptor_common(
-    oprt_operator_descriptor_t *desc,
-    const oprt_tensor_view_t *out,
-    const oprt_tensor_view_t *in,
-    float negative_slope) {
-    if (desc == nullptr) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    *desc = nullptr;
-
-    if (out == nullptr || in == nullptr) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-
-    auto *typed = new ReluDescriptor();
-    auto status = oprt::init_elementwise_descriptor(typed, out, {in});
-    if (status != OPRT_SUCCESS) {
-        delete typed;
-        return status;
-    }
-    typed->negative_slope = negative_slope;
-    *desc = typed;
-    return OPRT_SUCCESS;
-}
-
-inline oprt_status_t get_relu_workspace_size_common(oprt_operator_descriptor_t desc,
-                                                    size_t *size) {
-    return oprt::get_elementwise_workspace_size(desc, size);
-}
-
-inline oprt_status_t validate_relu_execute_args(oprt_operator_descriptor_t desc,
-                                                size_t workspace_size,
-                                                void *out,
-                                                const void *in) {
-    if (desc == nullptr || out == nullptr || in == nullptr) {
-        return OPRT_ERR_INVALID_ARG;
-    }
-    return oprt::validate_elementwise_execute_args(desc, workspace_size, out, {in});
-}
-
-inline oprt_status_t destroy_relu_descriptor_common(oprt_operator_descriptor_t desc) {
-    return oprt::destroy_elementwise_descriptor(desc);
-}
-
-} // namespace oprt::elementwise
--- a/ops/reduce_sum/metax/README.md
+++ b/ops/reduce_sum/metax/README.md
@ -1,4 +0,0 @@
-# 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.maca
+++ b/ops/reduce_sum/metax/reduce_sum_metax.maca
@ -1,132 +0,0 @@
-#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
@ -1,19 +0,0 @@
-#pragma once
-
-#include <cuda_runtime.h>
-#include <stdint.h>
-
-namespace oprt::reduce_sum::nvidia {
-
-__global__ void reduce_sum_rowwise_kernel(float *out, const float *in, int64_t rows, int64_t cols) {
-    // 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,105 +0,0 @@
-#include "operator_runtime/ops/reduce_sum.h"
-
-#include "operator_runtime/descriptor.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>
-
-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";
-    }
-};
-
-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;
-    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(
-    oprt_operator_descriptor_t desc) {
-    delete desc;
-    return OPRT_SUCCESS;
-}
--- a/ops/reduce_sum/tilelang/kernel.py
+++ b/ops/reduce_sum/tilelang/kernel.py
@ -1,29 +0,0 @@
-from __future__ import annotations
-
-import tilelang
-import tilelang.language as T
-
-
-@tilelang.jit(
-    pass_configs={
-        tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True,
-    },
-)
-def reduce_sum_kernel(src, BLOCK_N: int, BLOCK_M: int):
-    N, M = T.const("N, M")
-    dtype = T.float32
-    src: T.Tensor((N, M), dtype)
-    out = T.empty((N,), dtype)
-
-    # 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
@ -1,70 +0,0 @@
-from __future__ import annotations
-
-from dataclasses import dataclass
-from functools import lru_cache
-
-import torch
-
-from ops.reduce_sum.tilelang.kernel import reduce_sum_kernel
-
-@lru_cache(maxsize=32)
-def _compiled_reduce_sum(n: int, m: int, block_n: int, block_m: int):
-    return reduce_sum_kernel.compile(N=n, M=m, BLOCK_N=block_n, BLOCK_M=block_m)
-
-
-def _blocks(n: int, m: int) -> tuple[int, int]:
-    block_n = 16 if n % 16 == 0 else 1
-    block_m = 128 if m % 128 == 0 else m
-    return block_n, block_m
-
-
-@dataclass
-class TileLangReduceSumPrepared:
-    out: torch.Tensor
-    src: torch.Tensor
-    dim: int = 1
-    kernel: object | None = 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
-        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
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-
-def prepare_reduce_sum_tl(out: torch.Tensor, src: torch.Tensor, dim: int = 1) -> TileLangReduceSumPrepared:
-    if dim != 1:
-        raise ValueError("TileLang reduce_sum v1 supports dim=1")
-    n, m = src.shape
-    block_n, block_m = _blocks(n, m)
-    if n % block_n != 0 or m % block_m != 0:
-        raise ValueError("TileLang reduce_sum v1 requires divisible block sizes")
-    kernel = _compiled_reduce_sum(n, m, block_n, block_m)
-    return TileLangReduceSumPrepared(out, src, dim, kernel)
-
-
-def reduce_sum_tl_(out: torch.Tensor, src: torch.Tensor, dim: int = 1) -> torch.Tensor:
-    with prepare_reduce_sum_tl(out, src, dim) as prepared:
-        prepared.run()
-    return out
-
-
-def reduce_sum_tl(src: torch.Tensor, dim: int = 1) -> torch.Tensor:
-    out = torch.empty((src.shape[0],), dtype=src.dtype, device=src.device)
-    return reduce_sum_tl_(out, src, dim)
--- a/ops/softmax/metax/README.md
+++ b/ops/softmax/metax/README.md
@ -1,4 +0,0 @@
-# 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.maca
+++ b/ops/softmax/metax/softmax_metax.maca
@ -1,155 +0,0 @@
-#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
@ -1,22 +0,0 @@
-#pragma once
-
-#include <cuda_runtime.h>
-#include <float.h>
-#include <math.h>
-#include <stdint.h>
-
-namespace oprt::softmax::nvidia {
-
-__global__ void softmax_rowwise_kernel(float *out, const float *in, int64_t rows, int64_t cols) {
-    // 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,105 +0,0 @@
-#include "operator_runtime/ops/softmax.h"
-
-#include "operator_runtime/descriptor.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>
-
-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";
-    }
-};
-
-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;
-    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(
-    oprt_operator_descriptor_t desc) {
-    delete desc;
-    return OPRT_SUCCESS;
-}
--- a/ops/softmax/tilelang/kernel.py
+++ b/ops/softmax/tilelang/kernel.py
@ -1,36 +0,0 @@
-from __future__ import annotations
-
-import tilelang
-import tilelang.language as T
-
-
-@tilelang.jit(
-    pass_configs={
-        tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True,
-    },
-)
-def softmax_kernel(src, BLOCK_N: int, BLOCK_M: int):
-    log2_e = 1.44269504
-    N, M = T.const("N, M")
-    dtype = T.float32
-    src: T.Tensor((N, M), dtype)
-    out = T.empty((N, M), dtype)
-
-    # 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
@ -1,70 +0,0 @@
-from __future__ import annotations
-
-from dataclasses import dataclass
-from functools import lru_cache
-
-import torch
-
-from ops.softmax.tilelang.kernel import softmax_kernel
-
-@lru_cache(maxsize=32)
-def _compiled_softmax(n: int, m: int, block_n: int, block_m: int):
-    return softmax_kernel.compile(N=n, M=m, BLOCK_N=block_n, BLOCK_M=block_m)
-
-
-def _blocks(n: int, m: int) -> tuple[int, int]:
-    block_n = 16 if n % 16 == 0 else 1
-    block_m = 256 if m % 256 == 0 else m
-    return block_n, block_m
-
-
-@dataclass
-class TileLangSoftmaxPrepared:
-    out: torch.Tensor
-    src: torch.Tensor
-    dim: int = 1
-    kernel: object | None = 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
-        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
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-
-def prepare_softmax_tl(out: torch.Tensor, src: torch.Tensor, dim: int = 1) -> TileLangSoftmaxPrepared:
-    if dim != 1:
-        raise ValueError("TileLang softmax v1 supports dim=1")
-    n, m = src.shape
-    block_n, block_m = _blocks(n, m)
-    if n % block_n != 0 or m % block_m != 0:
-        raise ValueError("TileLang softmax v1 requires divisible block sizes")
-    kernel = _compiled_softmax(n, m, block_n, block_m)
-    return TileLangSoftmaxPrepared(out, src, dim, kernel)
-
-
-def softmax_tl_(out: torch.Tensor, src: torch.Tensor, dim: int = 1) -> torch.Tensor:
-    with prepare_softmax_tl(out, src, dim) as prepared:
-        prepared.run()
-    return out
-
-
-def softmax_tl(src: torch.Tensor, dim: int = 1) -> torch.Tensor:
-    out = torch.empty_like(src)
-    return softmax_tl_(out, src, dim)
--- a/ops/vector_add/metax/README.md
+++ b/ops/vector_add/metax/README.md
@ -1,4 +0,0 @@
-# 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.maca
+++ b/ops/vector_add/metax/vector_add_metax.maca
@ -1,141 +0,0 @@
-#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
@ -1,34 +0,0 @@
-#pragma once
-
-#include <cuda_fp16.h>
-#include <stdint.h>
-
-namespace oprt::vector_add::nvidia {
-
-template <typename T>
-__device__ T add_values(T a, T b) {
-    // TODO: return the elementwise sum for generic types.
-    return T{};
-}
-
-template <>
-__device__ inline half add_values<half>(half a, half b) {
-    // TODO: return the half-precision elementwise sum.
-    return half{};
-}
-
-template <typename T>
-__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,119 +0,0 @@
-#include "operator_runtime/ops/vector_add.h"
-
-#include "operator_runtime/descriptor.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>
-
-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>
-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);
-    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;
-}
-
-} // 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/tilelang/kernel.py
+++ b/ops/vector_add/tilelang/kernel.py
@ -1,21 +0,0 @@
-from __future__ import annotations
-
-import tilelang
-import tilelang.language as T
-
-
-@tilelang.jit
-def vector_add_kernel(a, b, BLOCK_N: int, dtype):
-    N = T.const("N")
-    a: T.Tensor((N,), dtype)
-    b: T.Tensor((N,), dtype)
-    out = T.empty((N,), dtype)
-
-    # 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
@ -1,74 +0,0 @@
-from __future__ import annotations
-
-from dataclasses import dataclass
-from functools import lru_cache
-
-import tilelang.language as T
-import torch
-
-from ops.vector_add.tilelang.kernel import vector_add_kernel
-
-
-def _tl_dtype(dtype: torch.dtype):
-    if dtype is torch.float16:
-        return T.float16
-    if dtype is torch.float32:
-        return T.float32
-    raise TypeError(f"unsupported TileLang dtype: {dtype}")
-
-def _block_n(n: int) -> int:
-    return 1024 if n % 1024 == 0 else n
-
-
-@lru_cache(maxsize=32)
-def _compiled_vector_add(n: int, block_n: int, dtype: torch.dtype):
-    return vector_add_kernel.compile(N=n, BLOCK_N=block_n, dtype=_tl_dtype(dtype))
-
-
-@dataclass
-class TileLangVectorAddPrepared:
-    out: torch.Tensor
-    a: torch.Tensor
-    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.run_inputs())
-
-    def destroy(self) -> None:
-        pass
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-
-def prepare_vector_add_tl(out: torch.Tensor, a: torch.Tensor, b: torch.Tensor) -> TileLangVectorAddPrepared:
-    n = out.numel()
-    block_n = _block_n(n)
-    if n % block_n != 0:
-        raise ValueError("TileLang vector_add v1 requires N % BLOCK_N == 0")
-    kernel = _compiled_vector_add(n, block_n, out.dtype)
-    return TileLangVectorAddPrepared(out, a, b, kernel)
-
-
-def vector_add_tl_(out: torch.Tensor, a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
-    with prepare_vector_add_tl(out, a, b) as prepared:
-        prepared.run()
-    return out
-
-
-def vector_add_tl(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
-    out = torch.empty_like(a)
-    return vector_add_tl_(out, a, b)
--- a/python/operator_runtime/init.py
+++ b/python/operator_runtime/init.py
@ -1,18 +0,0 @@
-from .backend import Backend, normalize_backend
-from .ops import (
-    copy, copy_, prepare_copy,
-    vector_add, vector_add_, prepare_vector_add,
-    reduce_sum, reduce_sum_, prepare_reduce_sum,
-    relu, relu_, prepare_relu,
-    softmax, softmax_, prepare_softmax,
-)
-
-__all__ = [
-    "Backend",
-    "normalize_backend",
-    "copy", "copy_", "prepare_copy",
-    "vector_add", "vector_add_", "prepare_vector_add",
-    "reduce_sum", "reduce_sum_", "prepare_reduce_sum",
-    "relu", "relu_", "prepare_relu",
-    "softmax", "softmax_", "prepare_softmax",
-]
--- a/python/operator_runtime/_internal/init.py
+++ b/python/operator_runtime/_internal/init.py
@ -1,35 +0,0 @@
-from .loader import load_library
-from .tensor_view import TensorView, tensor_view, dtype_to_oprt, current_stream_ptr, OPRT_MAX_DIMS
-from .bindings import (
-    CFunctions,
-    Descriptor,
-    OperatorRuntimeError,
-    Status,
-    bind_elementwise,
-    bind_unary,
-    bind_relu,
-    bind_binary,
-    bind_reduce_like,
-    check_status,
-)
-from .prepared import PreparedOp
-
-__all__ = [
-    "load_library",
-    "TensorView",
-    "tensor_view",
-    "dtype_to_oprt",
-    "current_stream_ptr",
-    "OPRT_MAX_DIMS",
-    "CFunctions",
-    "Descriptor",
-    "OperatorRuntimeError",
-    "Status",
-    "bind_elementwise",
-    "bind_unary",
-    "bind_relu",
-    "bind_binary",
-    "bind_reduce_like",
-    "check_status",
-    "PreparedOp",
-]
--- a/python/operator_runtime/_internal/bindings.py
+++ b/python/operator_runtime/_internal/bindings.py
@ -1,140 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-from dataclasses import dataclass
-from typing import Callable, Sequence
-
-from operator_runtime.backend import Backend
-from operator_runtime.backend import normalize_backend
-
-from .loader import load_library
-from .tensor_view import TensorView
-
-Status = ctypes.c_int
-Descriptor = ctypes.c_void_p
-OPRT_BACKEND_NVIDIA = 1
-OPRT_BACKEND_METAX = 2
-
-
-class OperatorRuntimeError(RuntimeError):
-    pass
-
-
-def _backend_code(backend: str | Backend) -> int:
-    normalized = normalize_backend(backend)
-    if normalized is Backend.NVIDIA:
-        return OPRT_BACKEND_NVIDIA
-    if normalized is Backend.METAX:
-        return OPRT_BACKEND_METAX
-    raise NotImplementedError(f"backend {normalized.value} is not runnable through the C ABI")
-
-
-def check_status(status: int, lib=None) -> None:
-    if status == 0:
-        return
-    if lib is None:
-        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)
-
-
-def set_runtime_backend(lib, backend: str | Backend) -> None:
-    lib.oprt_set_backend.argtypes = [ctypes.c_int]
-    lib.oprt_set_backend.restype = Status
-    check_status(lib.oprt_set_backend(_backend_code(backend)), lib)
-
-
-@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_lifecycle(
-    name: str,
-    backend: str | Backend,
-    create_argtypes: Sequence[object],
-    execute_argtypes: Sequence[object],
-) -> CFunctions:
-    lib = load_library(backend)
-    set_runtime_backend(lib, backend)
-    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), *create_argtypes]
-    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, *execute_argtypes, 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_elementwise(
-    name: str,
-    input_count: int,
-    scalar_argtypes: Sequence[object] = (),
-    backend: str | Backend = Backend.NVIDIA,
-) -> CFunctions:
-    if input_count <= 0:
-        raise ValueError("elementwise operators require at least one input")
-    tensor_views = [ctypes.POINTER(TensorView)] * (input_count + 1)
-    data_ptrs = [ctypes.c_void_p] * (input_count + 1)
-    return _bind_lifecycle(name, backend, [*tensor_views, *scalar_argtypes], data_ptrs)
-
-
-def bind_unary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions:
-    return bind_elementwise(name, 1, backend=backend)
-
-
-def bind_relu(backend: str | Backend = Backend.NVIDIA) -> CFunctions:
-    return bind_elementwise("relu", 1, [ctypes.c_float], backend)
-
-
-def bind_binary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions:
-    return bind_elementwise(name, 2, backend=backend)
-
-
-def bind_reduce_like(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions:
-    return _bind_lifecycle(
-        name,
-        backend,
-        [
-            ctypes.POINTER(TensorView),
-            ctypes.POINTER(TensorView),
-            ctypes.c_int64,
-        ],
-        [ctypes.c_void_p, ctypes.c_void_p],
-    )
--- a/python/operator_runtime/_internal/loader.py
+++ b/python/operator_runtime/_internal/loader.py
@ -1,72 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-import os
-from functools import lru_cache
-from pathlib import Path
-
-
-def _backend_key(backend: object | None) -> str | None:
-    if backend is None:
-        return None
-    from operator_runtime.backend import normalize_backend
-
-    return normalize_backend(backend).value
-
-
-def _append_build_dir(candidates: list[Path], build_dir: str | None) -> None:
-    if not build_dir:
-        return
-    root = Path(build_dir)
-    candidates.extend([root / "libcamp_ops.so", root / "ops" / "libcamp_ops.so"])
-
-
-def _candidate_library_paths(backend: str | None = None) -> list[Path]:
-    candidates: list[Path] = []
-    repo_root = Path(__file__).resolve().parents[3]
-
-    if backend == "nvidia":
-        _append_build_dir(candidates, os.environ.get("CAMP_NVIDIA_BUILD_DIR"))
-        _append_build_dir(candidates, os.environ.get("CAMP_BUILD_DIR"))
-        candidates.extend(
-            [
-                repo_root / "build" / "libcamp_ops.so",
-                repo_root / "build" / "ops" / "libcamp_ops.so",
-            ]
-        )
-    elif backend == "metax":
-        _append_build_dir(candidates, os.environ.get("CAMP_METAX_BUILD_DIR"))
-        _append_build_dir(candidates, os.environ.get("CAMP_BUILD_DIR"))
-        candidates.extend(
-            [
-                repo_root / "build-metax" / "libcamp_ops.so",
-                repo_root / "build-metax" / "ops" / "libcamp_ops.so",
-            ]
-        )
-    else:
-        _append_build_dir(candidates, os.environ.get("CAMP_BUILD_DIR"))
-        _append_build_dir(candidates, os.environ.get("CAMP_NVIDIA_BUILD_DIR"))
-        _append_build_dir(candidates, os.environ.get("CAMP_METAX_BUILD_DIR"))
-        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
-
-
-@lru_cache(maxsize=None)
-def _load_library(backend: str | None) -> ctypes.CDLL:
-    for path in _candidate_library_paths(backend):
-        if path.exists():
-            return ctypes.CDLL(str(path))
-    searched = ", ".join(str(path) for path in _candidate_library_paths(backend))
-    label = "default" if backend is None else backend
-    raise FileNotFoundError(f"libcamp_ops.so for backend {label} not found; searched: {searched}")
-
-
-def load_library(backend: object | None = None) -> ctypes.CDLL:
-    return _load_library(_backend_key(backend))
--- a/python/operator_runtime/_internal/prepared.py
+++ b/python/operator_runtime/_internal/prepared.py
@ -1,65 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-from dataclasses import dataclass
-from typing import Any
-
-import torch
-
-from .bindings import CFunctions, Descriptor, check_status
-
-
-@dataclass
-class PreparedOp:
-    funcs: CFunctions
-    descriptor: Descriptor
-    workspace: torch.Tensor | None
-    runner_args: tuple[Any, ...]
-    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()
-        status = self.funcs.execute(
-            self.descriptor,
-            workspace_ptr,
-            workspace_size,
-            *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))
-            self.descriptor = Descriptor()
-
-    def __enter__(self) -> "PreparedOp":
-        return self
-
-    def __exit__(self, exc_type, exc, tb) -> None:
-        self.destroy()
-
-    def __del__(self) -> None:
-        try:
-            self.destroy()
-        except Exception:
-            pass
--- a/python/operator_runtime/_internal/tensor_view.py
+++ b/python/operator_runtime/_internal/tensor_view.py
@ -1,54 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-
-import torch
-
-OPRT_MAX_DIMS = 8
-
-OPRT_DTYPE_F16 = 0
-OPRT_DTYPE_F32 = 1
-
-
-class TensorView(ctypes.Structure):
-    _fields_ = [
-        ("data", ctypes.c_void_p),
-        ("dtype", ctypes.c_int),
-        ("ndim", ctypes.c_int32),
-        ("shape", ctypes.c_int64 * OPRT_MAX_DIMS),
-        ("strides", ctypes.c_int64 * OPRT_MAX_DIMS),
-    ]
-
-
-def dtype_to_oprt(dtype: torch.dtype) -> int:
-    if dtype is torch.float16:
-        return OPRT_DTYPE_F16
-    if dtype is torch.float32:
-        return OPRT_DTYPE_F32
-    raise TypeError(f"unsupported dtype: {dtype}")
-
-
-def tensor_view(tensor: torch.Tensor) -> TensorView:
-    if tensor.ndim > OPRT_MAX_DIMS:
-        raise ValueError(f"ndim {tensor.ndim} exceeds OPRT_MAX_DIMS={OPRT_MAX_DIMS}")
-    shape = (ctypes.c_int64 * OPRT_MAX_DIMS)()
-    strides = (ctypes.c_int64 * OPRT_MAX_DIMS)()
-    for i, dim in enumerate(tensor.shape):
-        shape[i] = dim
-    for i, stride in enumerate(tensor.stride()):
-        strides[i] = stride
-    return TensorView(
-        ctypes.c_void_p(tensor.data_ptr()),
-        dtype_to_oprt(tensor.dtype),
-        tensor.ndim,
-        shape,
-        strides,
-    )
-
-
-def current_stream_ptr(tensor: torch.Tensor | None = None) -> ctypes.c_void_p:
-    if not torch.cuda.is_available():
-        return 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/backend.py
+++ b/python/operator_runtime/backend.py
@ -1,19 +0,0 @@
-from __future__ import annotations
-
-from enum import Enum
-
-
-class Backend(str, Enum):
-    NVIDIA = "nvidia"
-    TILELANG = "tilelang"
-    METAX = "metax"
-
-
-def normalize_backend(backend: str | Backend) -> Backend:
-    if isinstance(backend, Backend):
-        return backend
-    try:
-        return Backend(backend.lower())
-    except ValueError as exc:
-        raise ValueError(f"unknown backend: {backend}") from exc
-
--- a/python/operator_runtime/ops/init.py
+++ b/python/operator_runtime/ops/init.py
@ -1,13 +0,0 @@
-from .copy import copy, copy_, prepare_copy
-from .vector_add import vector_add, vector_add_, prepare_vector_add
-from .reduce_sum import reduce_sum, reduce_sum_, prepare_reduce_sum
-from .relu import relu, relu_, prepare_relu
-from .softmax import softmax, softmax_, prepare_softmax
-
-__all__ = [
-    "copy", "copy_", "prepare_copy",
-    "vector_add", "vector_add_", "prepare_vector_add",
-    "reduce_sum", "reduce_sum_", "prepare_reduce_sum",
-    "relu", "relu_", "prepare_relu",
-    "softmax", "softmax_", "prepare_softmax",
-]
--- a/python/operator_runtime/ops/_common.py
+++ b/python/operator_runtime/ops/_common.py
@ -1,99 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-from dataclasses import dataclass
-from typing import Callable, Sequence
-
-import torch
-
-from operator_runtime.backend import Backend, normalize_backend
-from operator_runtime._internal import CFunctions, Descriptor, PreparedOp, bind_elementwise, check_status, tensor_view
-
-
-def _broadcastable_to(src_shape: tuple[int, ...], out_shape: tuple[int, ...]) -> bool:
-    if len(src_shape) > len(out_shape):
-        return False
-    offset = len(out_shape) - len(src_shape)
-    for out_dim, out_extent in enumerate(out_shape):
-        src_dim = out_dim - offset
-        if src_dim < 0:
-            continue
-        src_extent = src_shape[src_dim]
-        if src_extent != out_extent and src_extent != 1:
-            return False
-    return True
-
-
-def _has_broadcast_dim(tensor: torch.Tensor) -> bool:
-    return any(size > 1 and stride == 0 for size, stride in zip(tensor.shape, tensor.stride()))
-
-
-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)
-
-
-@dataclass(frozen=True)
-class ElementwiseOpSpec:
-    name: str
-    input_count: int
-    scalar_argtypes: tuple[object, ...] = ()
-    allow_broadcast: bool = True
-
-
-def check_elementwise_tensors(spec: ElementwiseOpSpec, out: torch.Tensor, inputs: Sequence[torch.Tensor]) -> None:
-    if len(inputs) != spec.input_count:
-        raise ValueError(f"{spec.name} expects {spec.input_count} inputs, got {len(inputs)}")
-    if not out.is_cuda or any(not tensor.is_cuda for tensor in inputs):
-        raise ValueError(f"{spec.name} expects CUDA tensors")
-    if _has_broadcast_dim(out):
-        raise ValueError(f"{spec.name} expects a writable output tensor")
-    for tensor in inputs:
-        if spec.allow_broadcast:
-            if not _broadcastable_to(tuple(tensor.shape), tuple(out.shape)):
-                raise ValueError(f"{spec.name} expects inputs to be broadcastable to out")
-        elif out.shape != tensor.shape:
-            raise ValueError(f"{spec.name} expects matching shapes")
-        if out.dtype != tensor.dtype:
-            raise TypeError(f"{spec.name} expects matching dtypes")
-
-
-def prepare_elementwise_op(
-    spec: ElementwiseOpSpec,
-    out: torch.Tensor,
-    inputs: Sequence[torch.Tensor],
-    scalars: Sequence[object] = (),
-    backend: str | Backend = Backend.NVIDIA,
-    scalar_converters: Sequence[Callable[[object], object]] = (),
-) -> PreparedOp:
-    backend = normalize_backend(backend)
-    if backend not in (Backend.NVIDIA, Backend.METAX):
-        raise NotImplementedError(f"backend {backend.value} is not runnable")
-
-    check_elementwise_tensors(spec, out, inputs)
-    if len(scalars) != len(spec.scalar_argtypes):
-        raise ValueError(f"{spec.name} expects {len(spec.scalar_argtypes)} scalar args, got {len(scalars)}")
-    if scalar_converters and len(scalar_converters) != len(scalars):
-        raise ValueError("scalar_converters must match scalars")
-
-    funcs = bind_elementwise(spec.name, spec.input_count, spec.scalar_argtypes, backend)
-    views = [tensor_view(out), *(tensor_view(tensor) for tensor in inputs)]
-    converters: Sequence[Callable[[object], object]] = scalar_converters or tuple(lambda value: value for _ in scalars)
-    converted_scalars = [
-        argtype(converter(value))
-        for value, argtype, converter in zip(scalars, spec.scalar_argtypes, converters)
-    ]
-    create_args = (*[ctypes.byref(view) for view in views], *converted_scalars)
-    return build_prepared_op(funcs, create_args, (out, *inputs), out)
--- a/python/operator_runtime/ops/copy.py
+++ b/python/operator_runtime/ops/copy.py
@ -1,48 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-
-import torch
-
-from operator_runtime.backend import Backend, normalize_backend
-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:
-    if not out.is_cuda or not src.is_cuda:
-        raise ValueError("copy expects CUDA tensors")
-    if out.shape != src.shape:
-        raise ValueError("copy expects matching shapes")
-    if out.dtype != src.dtype:
-        raise TypeError("copy expects matching dtypes")
-    if not out.is_contiguous() or not src.is_contiguous():
-        raise ValueError("copy v1 supports contiguous tensors only")
-
-
-def prepare_copy(out: torch.Tensor, src: torch.Tensor, backend: str | Backend = Backend.NVIDIA) -> PreparedOp:
-    backend = normalize_backend(backend)
-    _check(out, src)
-    if backend is Backend.TILELANG:
-        from ops.copy.tilelang.copy_tl import prepare_copy_tl
-
-        return prepare_copy_tl(out, src)
-    if backend not in (Backend.NVIDIA, Backend.METAX):
-        raise NotImplementedError(f"backend {backend.value} is not runnable")
-
-    funcs = bind_unary("copy", backend)
-    out_view = tensor_view(out)
-    src_view = tensor_view(src)
-    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:
-    with prepare_copy(out, src, backend) as prepared:
-        prepared.run()
-    return out
-
-
-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
@ -1,57 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-
-import torch
-
-from operator_runtime.backend import Backend, normalize_backend
-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:
-    if not out.is_cuda or not src.is_cuda:
-        raise ValueError("reduce_sum expects CUDA tensors")
-    if src.dtype is not torch.float32 or out.dtype is not torch.float32:
-        raise TypeError("reduce_sum v1 supports float32 only")
-    if src.ndim != 2 or out.ndim != 1 or dim != 1:
-        raise ValueError("reduce_sum v1 supports 2D row-wise reduction over dim=1")
-    if out.shape[0] != src.shape[0]:
-        raise ValueError("reduce_sum output shape must be [src.shape[0]]")
-    if not out.is_contiguous() or not src.is_contiguous():
-        raise ValueError("reduce_sum v1 supports contiguous tensors only")
-
-
-def prepare_reduce_sum(
-    out: torch.Tensor,
-    src: torch.Tensor,
-    dim: int = 1,
-    backend: str | Backend = Backend.NVIDIA,
-) -> PreparedOp:
-    backend = normalize_backend(backend)
-    _check(out, src, dim)
-    if backend is Backend.TILELANG:
-        from ops.reduce_sum.tilelang.reduce_sum_tl import prepare_reduce_sum_tl
-
-        return prepare_reduce_sum_tl(out, src, dim=dim)
-    if backend not in (Backend.NVIDIA, Backend.METAX):
-        raise NotImplementedError(f"backend {backend.value} is not runnable")
-
-    funcs = bind_reduce_like("reduce_sum", backend)
-    out_view = tensor_view(out)
-    src_view = tensor_view(src)
-    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:
-    with prepare_reduce_sum(out, src, dim, backend) as prepared:
-        prepared.run()
-    return out
-
-
-def reduce_sum(src: torch.Tensor, dim: int = 1, backend: str | Backend = Backend.NVIDIA) -> torch.Tensor:
-    if dim != 1 or src.ndim != 2:
-        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/relu.py
+++ b/python/operator_runtime/ops/relu.py
@ -1,52 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-
-import torch
-
-from operator_runtime.backend import Backend
-from operator_runtime._internal import PreparedOp
-from operator_runtime.ops._common import ElementwiseOpSpec, prepare_elementwise_op
-
-
-_RELU_SPEC = ElementwiseOpSpec(
-    name="relu",
-    input_count=1,
-    scalar_argtypes=(ctypes.c_float,),
-)
-
-
-def prepare_relu(
-    out: torch.Tensor,
-    src: torch.Tensor,
-    negative_slope: float = 0.0,
-    backend: str | Backend = Backend.NVIDIA,
-) -> PreparedOp:
-    return prepare_elementwise_op(
-        _RELU_SPEC,
-        out,
-        (src,),
-        (negative_slope,),
-        backend,
-        (float,),
-    )
-
-
-def relu_(
-    out: torch.Tensor,
-    src: torch.Tensor,
-    negative_slope: float = 0.0,
-    backend: str | Backend = Backend.NVIDIA,
-) -> torch.Tensor:
-    with prepare_relu(out, src, negative_slope, backend) as prepared:
-        prepared.run()
-    return out
-
-
-def relu(
-    src: torch.Tensor,
-    negative_slope: float = 0.0,
-    backend: str | Backend = Backend.NVIDIA,
-) -> torch.Tensor:
-    out = torch.empty_like(src)
-    return relu_(out, src, negative_slope, backend)
--- a/python/operator_runtime/ops/softmax.py
+++ b/python/operator_runtime/ops/softmax.py
@ -1,55 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-
-import torch
-
-from operator_runtime.backend import Backend, normalize_backend
-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:
-    if not out.is_cuda or not src.is_cuda:
-        raise ValueError("softmax expects CUDA tensors")
-    if src.dtype is not torch.float32 or out.dtype is not torch.float32:
-        raise TypeError("softmax v1 supports float32 only")
-    if src.ndim != 2 or out.ndim != 2 or dim != 1:
-        raise ValueError("softmax v1 supports 2D row-wise dim=1")
-    if out.shape != src.shape:
-        raise ValueError("softmax output shape must match input")
-    if not out.is_contiguous() or not src.is_contiguous():
-        raise ValueError("softmax v1 supports contiguous tensors only")
-
-
-def prepare_softmax(
-    out: torch.Tensor,
-    src: torch.Tensor,
-    dim: int = 1,
-    backend: str | Backend = Backend.NVIDIA,
-) -> PreparedOp:
-    backend = normalize_backend(backend)
-    _check(out, src, dim)
-    if backend is Backend.TILELANG:
-        from ops.softmax.tilelang.softmax_tl import prepare_softmax_tl
-
-        return prepare_softmax_tl(out, src, dim=dim)
-    if backend not in (Backend.NVIDIA, Backend.METAX):
-        raise NotImplementedError(f"backend {backend.value} is not runnable")
-
-    funcs = bind_reduce_like("softmax", backend)
-    out_view = tensor_view(out)
-    src_view = tensor_view(src)
-    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:
-    with prepare_softmax(out, src, dim, backend) as prepared:
-        prepared.run()
-    return out
-
-
-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
@ -1,54 +0,0 @@
-from __future__ import annotations
-
-import ctypes
-
-import torch
-
-from operator_runtime.backend import Backend, normalize_backend
-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:
-    if not out.is_cuda or not a.is_cuda or not b.is_cuda:
-        raise ValueError("vector_add expects CUDA tensors")
-    if out.shape != a.shape or out.shape != b.shape:
-        raise ValueError("vector_add v1 expects matching shapes")
-    if out.dtype != a.dtype or out.dtype != b.dtype:
-        raise TypeError("vector_add expects matching dtypes")
-    if not out.is_contiguous() or not a.is_contiguous() or not b.is_contiguous():
-        raise ValueError("vector_add v1 supports contiguous tensors only")
-
-
-def prepare_vector_add(
-    out: torch.Tensor,
-    a: torch.Tensor,
-    b: torch.Tensor,
-    backend: str | Backend = Backend.NVIDIA,
-) -> PreparedOp:
-    backend = normalize_backend(backend)
-    _check(out, a, b)
-    if backend is Backend.TILELANG:
-        from ops.vector_add.tilelang.vector_add_tl import prepare_vector_add_tl
-
-        return prepare_vector_add_tl(out, a, b)
-    if backend not in (Backend.NVIDIA, Backend.METAX):
-        raise NotImplementedError(f"backend {backend.value} is not runnable")
-
-    funcs = bind_binary("vector_add", backend)
-    out_view = tensor_view(out)
-    a_view = tensor_view(a)
-    b_view = tensor_view(b)
-    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:
-    with prepare_vector_add(out, a, b, backend) as prepared:
-        prepared.run()
-    return out
-
-
-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/init.py
+++ b/python/operator_runtime_testing/init.py
@ -1,10 +0,0 @@
-from .assertions import require_cuda, assert_close
-from .benchmark import cuda_time_ms
-from .profiler import PerformanceResult
-
-__all__ = [
-    "require_cuda",
-    "assert_close",
-    "cuda_time_ms",
-    "PerformanceResult",
-]
--- a/python/operator_runtime_testing/assertions.py
+++ b/python/operator_runtime_testing/assertions.py
@ -1,13 +0,0 @@
-from __future__ import annotations
-
-import pytest
-import torch
-
-
-def require_cuda() -> None:
-    if not torch.cuda.is_available():
-        pytest.skip("CUDA is required")
-
-
-def assert_close(actual: torch.Tensor, expected: torch.Tensor, *, atol: float, rtol: float) -> None:
-    torch.testing.assert_close(actual, expected, atol=atol, rtol=rtol)
--- a/python/operator_runtime_testing/benchmark.py
+++ b/python/operator_runtime_testing/benchmark.py
@ -1,136 +0,0 @@
-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
@ -1,19 +0,0 @@
-from __future__ import annotations
-
-from dataclasses import dataclass
-
-
-@dataclass
-class PerformanceResult:
-    operator: str
-    backend: str
-    shape: str
-    dtype: str
-    runtime_ms: float
-    torch_ms: float | None = None
-
-    @property
-    def speedup(self) -> float | None:
-        if self.torch_ms is None or self.runtime_ms <= 0:
-            return None
-        return self.torch_ms / self.runtime_ms
--- a/requirements.txt
+++ b/requirements.txt
@ -1,14 +0,0 @@
-# Python dependencies for operator_runtime_training
-
-# Runtime and benchmarking (required)
-torch>=2.0
-
-# TileLang backend (optional)
-tilelang
-
-# Build tools
-cmake>=3.22
-ninja
-
-# Testing framework
-pytest>=7.0
--- a/scripts/build_metax.sh
+++ b/scripts/build_metax.sh
@ -1,131 +0,0 @@
-#!/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}"
-CMAKE_GENERATOR="${CMAKE_GENERATOR:-Ninja}"
-CAMP_FORCE_RECONFIGURE="${CAMP_FORCE_RECONFIGURE:-0}"
-CAMP_USE_TILELANG_METAX="${CAMP_USE_TILELANG_METAX:-0}"
-CAMP_TILELANG_SOURCE_ROOT="${CAMP_TILELANG_SOURCE_ROOT:-/root/tilelang-metax}"
-
-MODE="${1:-build}"
-
-export MACA_PATH
-export CUCC_PATH
-export CUDA_PATH
-export LD_LIBRARY_PATH="${MACA_PATH}/lib:${LD_LIBRARY_PATH:-}"
-
-tilelang_env() {
-  if [[ "${CAMP_USE_TILELANG_METAX}" != "1" ]]; then
-    return 0
-  fi
-  if [[ ! -d "${CAMP_TILELANG_SOURCE_ROOT}" ]]; then
-    echo "TileLang source root not found: ${CAMP_TILELANG_SOURCE_ROOT}" >&2
-    return 1
-  fi
-  if [[ ! -d "${CAMP_TILELANG_SOURCE_ROOT}/build/lib" ]]; then
-    echo "TileLang source build output not found: ${CAMP_TILELANG_SOURCE_ROOT}/build/lib" >&2
-    echo "Build /root/tilelang-metax first with USE_MACA=ON cmake -S ... -B ... && make -C build" >&2
-    return 1
-  fi
-  export PYTHONPATH="${CAMP_TILELANG_SOURCE_ROOT}:${ROOT}/python:${ROOT}${PYTHONPATH:+:${PYTHONPATH}}"
-  export LD_LIBRARY_PATH="${CAMP_TILELANG_SOURCE_ROOT}/build/lib:${MACA_PATH}/lib:${LD_LIBRARY_PATH:-}"
-}
-
-prepare_build_dir() {
-  mkdir -p "${BUILD_DIR}"
-  if [[ "${CAMP_FORCE_RECONFIGURE}" == "1" ]]; then
-    rm -f "${BUILD_DIR}/CMakeCache.txt"
-    rm -rf "${BUILD_DIR}/CMakeFiles"
-  fi
-}
-
-configure() {
-  prepare_build_dir
-  "${PYTHON_BIN}" -m cmake -G "${CMAKE_GENERATOR}" "${ROOT}" \
-    -B "${BUILD_DIR}" \
-    -DCAMP_ENABLE_NVIDIA=OFF \
-    -DCAMP_ENABLE_METAX=ON
-}
-
-build() {
-  "${PYTHON_BIN}" -m cmake --build "${BUILD_DIR}" -- -v
-}
-
-test_pytest() {
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" -m pytest "${ROOT}/tests" -v --backend metax
-}
-
-test_run_ops() {
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/tests/run_ops.py" --op all --backend metax --mode all
-}
-
-test_examples() {
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/01_copy.py" --backend metax
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/02_vector_add.py" --backend metax
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/03_reduce_sum.py" --backend metax
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/04_softmax.py" --backend metax
-}
-
-test_tilelang() {
-  tilelang_env
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" -m pytest "${ROOT}/tests/op_tests/test_copy.py" -v --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" -m pytest "${ROOT}/tests/op_tests/test_vector_add.py" -v --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" -m pytest "${ROOT}/tests/op_tests/test_reduce_sum.py" -v --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" -m pytest "${ROOT}/tests/op_tests/test_softmax.py" -v --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/01_copy.py" --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/02_vector_add.py" --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/03_reduce_sum.py" --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/04_softmax.py" --backend tilelang
-  env CAMP_BUILD_DIR="${BUILD_DIR}" "${PYTHON_BIN}" "${ROOT}/examples/05_tilelang_copy_modes.py"
-}
-
-case "${MODE}" in
-  env)
-    echo "BUILD_DIR=${BUILD_DIR}"
-    echo "PYTHON_BIN=${PYTHON_BIN}"
-    echo "CMAKE_GENERATOR=${CMAKE_GENERATOR}"
-    echo "CAMP_FORCE_RECONFIGURE=${CAMP_FORCE_RECONFIGURE}"
-    echo "MACA_PATH=${MACA_PATH}"
-    echo "CUCC_PATH=${CUCC_PATH}"
-    echo "CUDA_PATH=${CUDA_PATH}"
-    echo "CAMP_USE_TILELANG_METAX=${CAMP_USE_TILELANG_METAX}"
-    echo "CAMP_TILELANG_SOURCE_ROOT=${CAMP_TILELANG_SOURCE_ROOT}"
-    ;;
-  configure)
-    configure
-    ;;
-  build)
-    configure
-    build
-    ;;
-  test)
-    test_pytest
-    test_run_ops
-    test_examples
-    if [[ "${CAMP_USE_TILELANG_METAX}" == "1" ]]; then
-      test_tilelang
-    fi
-    ;;
-  all)
-    configure
-    build
-    test_pytest
-    test_run_ops
-    test_examples
-    if [[ "${CAMP_USE_TILELANG_METAX}" == "1" ]]; then
-      test_tilelang
-    fi
-    ;;
-  clean)
-    rm -rf "${BUILD_DIR}"
-    ;;
-  *)
-    echo "usage: $0 {env|configure|build|test|all|clean}" >&2
-    exit 2
-    ;;
-esac
--- a/scripts/build_nvidia.sh
+++ b/scripts/build_nvidia.sh
@ -1,132 +0,0 @@
-#!/usr/bin/env bash
-set -euo pipefail
-
-ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
-BUILD_DIR="${BUILD_DIR:-${ROOT}/build-nvidia}"
-PYTHON_BIN="${PYTHON_BIN:-python3}"
-CMAKE_GENERATOR="${CMAKE_GENERATOR:-Ninja}"
-CAMP_ENABLE_CUTE="${CAMP_ENABLE_CUTE:-AUTO}"
-CMAKE_CUDA_ARCHITECTURES="${CMAKE_CUDA_ARCHITECTURES:-native}"
-CAMP_FORCE_RECONFIGURE="${CAMP_FORCE_RECONFIGURE:-0}"
-
-MODE="${1:-build}"
-
-CUTLASS_VERSION="${CUTLASS_VERSION:-v3.7.0}"
-CUTLASS_REPO="${CUTLASS_REPO:-https://github.com/NVIDIA/cutlass.git}"
-
-ensure_third_party() {
-  local tp_dir="${ROOT}/third_party"
-  if [[ -n "${CAMP_CUTLASS_ROOT:-}" ]]; then
-    return
-  fi
-  if [[ -f "${tp_dir}/cutlass/include/cute/tensor.hpp" ]]; then
-    CAMP_CUTLASS_ROOT="${tp_dir}/cutlass"
-    return
-  fi
-  mkdir -p "${tp_dir}"
-  echo "==> Fetching CUTLASS ${CUTLASS_VERSION} into third_party/cutlass ..."
-  git clone --depth 1 --branch "${CUTLASS_VERSION}" "${CUTLASS_REPO}" "${tp_dir}/cutlass"
-  CAMP_CUTLASS_ROOT="${tp_dir}/cutlass"
-}
-
-ensure_third_party
-
-if [[ -z "${CAMP_CUTLASS_ROOT:-}" && -n "${CUTLASS_ROOT:-}" ]]; then
-  CAMP_CUTLASS_ROOT="${CUTLASS_ROOT}"
-fi
-
-export CAMP_CUTLASS_ROOT="${CAMP_CUTLASS_ROOT:-}"
-
-prepare_build_dir() {
-  mkdir -p "${BUILD_DIR}"
-  if [[ "${CAMP_FORCE_RECONFIGURE}" == "1" ]]; then
-    rm -f "${BUILD_DIR}/CMakeCache.txt"
-    rm -rf "${BUILD_DIR}/CMakeFiles"
-  fi
-}
-
-configure() {
-  prepare_build_dir
-  local cmake_args=(
-    -G "${CMAKE_GENERATOR}"
-    "${ROOT}"
-    -B "${BUILD_DIR}"
-    -DCAMP_ENABLE_NVIDIA=ON
-    -DCAMP_ENABLE_METAX=OFF
-    -DCAMP_ENABLE_CUTE="${CAMP_ENABLE_CUTE}"
-    -DCMAKE_CUDA_ARCHITECTURES="${CMAKE_CUDA_ARCHITECTURES}"
-  )
-  if [[ -n "${CAMP_CUTLASS_ROOT}" ]]; then
-    cmake_args+=(-DCAMP_CUTLASS_ROOT="${CAMP_CUTLASS_ROOT}")
-  fi
-  "${PYTHON_BIN}" -m cmake "${cmake_args[@]}"
-}
-
-build() {
-  "${PYTHON_BIN}" -m cmake --build "${BUILD_DIR}" -- -v
-}
-
-test_pytest() {
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" -m pytest "${ROOT}/tests" -v --backend nvidia
-}
-
-test_run_ops() {
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" "${ROOT}/tests/run_ops.py" --op all --backend nvidia --mode all
-}
-
-test_examples() {
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" "${ROOT}/examples/01_copy.py" --backend nvidia
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" "${ROOT}/examples/02_vector_add.py" --backend nvidia
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" "${ROOT}/examples/03_reduce_sum.py" --backend nvidia
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" "${ROOT}/examples/04_softmax.py" --backend nvidia
-  env PYTHONPATH="${ROOT}/python:${ROOT}" CAMP_BUILD_DIR="${BUILD_DIR}" \
-    "${PYTHON_BIN}" "${ROOT}/examples/06_relu.py" --backend nvidia
-}
-
-show_env() {
-  echo "BUILD_DIR=${BUILD_DIR}"
-  echo "PYTHON_BIN=${PYTHON_BIN}"
-  echo "CMAKE_GENERATOR=${CMAKE_GENERATOR}"
-  echo "CAMP_ENABLE_CUTE=${CAMP_ENABLE_CUTE}"
-  echo "CMAKE_CUDA_ARCHITECTURES=${CMAKE_CUDA_ARCHITECTURES}"
-  echo "CAMP_FORCE_RECONFIGURE=${CAMP_FORCE_RECONFIGURE}"
-  echo "CAMP_CUTLASS_ROOT=${CAMP_CUTLASS_ROOT:-}"
-}
-
-case "${MODE}" in
-  env)
-    show_env
-    ;;
-  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 {env|configure|build|test|all|clean}" >&2
-    exit 2
-    ;;
-esac
--- a/tests/init.py
+++ b/tests/init.py
@ -1 +0,0 @@
-
--- a/tests/bench/init.py
+++ b/tests/bench/init.py
@ -1 +0,0 @@
-from __future__ import annotations
--- a/tests/bench/copy.py
+++ b/tests/bench/copy.py
@ -1,38 +0,0 @@
-from __future__ import annotations
-
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[2]
-PYTHON_DIR = ROOT / "python"
-if str(PYTHON_DIR) not in sys.path:
-    sys.path.insert(0, str(PYTHON_DIR))
-if str(ROOT) not in sys.path:
-    sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime_testing import cuda_time_ms, PerformanceResult
-from tests.cases import copy as copy_cases
-
-
-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)
-        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),
-                runtime,
-                torch_ms,
-            )
-        )
-    return rows
--- a/tests/bench/reduce_sum.py
+++ b/tests/bench/reduce_sum.py
@ -1,38 +0,0 @@
-from __future__ import annotations
-
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[2]
-PYTHON_DIR = ROOT / "python"
-if str(PYTHON_DIR) not in sys.path:
-    sys.path.insert(0, str(PYTHON_DIR))
-if str(ROOT) not in sys.path:
-    sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime_testing import cuda_time_ms, PerformanceResult
-from tests.cases import reduce_sum as reduce_sum_cases
-
-
-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")
-        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),
-                runtime,
-                torch_ms,
-            )
-        )
-    return rows
--- a/tests/bench/relu.py
+++ b/tests/bench/relu.py
@ -1,45 +0,0 @@
-from __future__ import annotations
-
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[2]
-PYTHON_DIR = ROOT / "python"
-if str(PYTHON_DIR) not in sys.path:
-    sys.path.insert(0, str(PYTHON_DIR))
-if str(ROOT) not in sys.path:
-    sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime_testing import cuda_time_ms, PerformanceResult
-from tests.cases import relu as relu_cases
-
-
-def bench_relu(backend: str) -> list[PerformanceResult]:
-    if backend not in ("nvidia", "metax"):
-        raise NotImplementedError("relu benchmark currently targets the C ABI elementwise framework")
-
-    rows: list[PerformanceResult] = []
-    for case in relu_cases.benchmark_cases():
-        src = torch.randn(case["shape"], dtype=case["dtype"], device="cuda")
-        out = torch.empty_like(src)
-        negative_slope = case["negative_slope"]
-        from operator_runtime import prepare_relu
-        with prepare_relu(out, src, negative_slope=negative_slope, backend=backend) as prepared:
-            runtime = cuda_time_ms(prepared.run_inputs, args=(src,))
-        torch_ms = cuda_time_ms(
-            lambda src: torch.where(src > 0, src, src * negative_slope, out=out),
-            args=(src,),
-        )
-        rows.append(
-            PerformanceResult(
-                "relu",
-                backend,
-                str(tuple(src.shape)),
-                str(src.dtype),
-                runtime,
-                torch_ms,
-            )
-        )
-    return rows
--- a/tests/bench/softmax.py
+++ b/tests/bench/softmax.py
@ -1,38 +0,0 @@
-from __future__ import annotations
-
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[2]
-PYTHON_DIR = ROOT / "python"
-if str(PYTHON_DIR) not in sys.path:
-    sys.path.insert(0, str(PYTHON_DIR))
-if str(ROOT) not in sys.path:
-    sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime_testing import cuda_time_ms, PerformanceResult
-from tests.cases import softmax as softmax_cases
-
-
-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")
-        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),
-                runtime,
-                torch_ms,
-            )
-        )
-    return rows
--- a/tests/bench/vector_add.py
+++ b/tests/bench/vector_add.py
@ -1,39 +0,0 @@
-from __future__ import annotations
-
-import sys
-from pathlib import Path
-
-ROOT = Path(__file__).resolve().parents[2]
-PYTHON_DIR = ROOT / "python"
-if str(PYTHON_DIR) not in sys.path:
-    sys.path.insert(0, str(PYTHON_DIR))
-if str(ROOT) not in sys.path:
-    sys.path.insert(0, str(ROOT))
-
-import torch
-
-from operator_runtime_testing import cuda_time_ms, PerformanceResult
-from tests.cases import vector_add as vector_add_cases
-
-
-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)
-        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),
-                runtime,
-                torch_ms,
-            )
-        )
-    return rows
--- a/tests/cases/init.py
+++ b/tests/cases/init.py
@ -1 +0,0 @@
-
--- a/tests/cases/copy.py
+++ b/tests/cases/copy.py
@ -1,31 +0,0 @@
-from __future__ import annotations
-
-import torch
-
-
-def correctness_cases():
-    return [
-        {"name": "contiguous_1k_fp32", "shape": (1024,), "dtype": torch.float32, "atol": 0, "rtol": 0},
-        {"name": "contiguous_1k_fp16", "shape": (1024,), "dtype": torch.float16, "atol": 0, "rtol": 0},
-        {"name": "contiguous_1536_fp32", "shape": (1536,), "dtype": torch.float32, "atol": 0, "rtol": 0},
-        {"name": "contiguous_1536_fp16", "shape": (1536,), "dtype": torch.float16, "atol": 0, "rtol": 0},
-    ]
-
-
-def api_error_cases():
-    return [
-        {"name": "shape_mismatch", "shape": (16,), "out_shape": (8,), "dtype": torch.float32},
-        {"name": "dtype_mismatch", "shape": (16,), "dtype": torch.float32, "out_dtype": torch.float16},
-        {"name": "non_contiguous", "shape": (4, 4), "dtype": torch.float32},
-    ]
-
-
-def benchmark_cases():
-    return [
-        {"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
@ -1,33 +0,0 @@
-from __future__ import annotations
-
-import torch
-
-
-def correctness_cases():
-    return [
-        {"name": "rowwise_1x128", "shape": (1, 128), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-        {"name": "rowwise_16x128", "shape": (16, 128), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-        {"name": "rowwise_16x256", "shape": (16, 256), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-        {"name": "rowwise_32x128", "shape": (32, 128), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-    ]
-
-
-def api_error_cases():
-    return [
-        {"name": "wrong_dim", "shape": (16, 16), "dtype": torch.float32, "dim": 0},
-        {"name": "wrong_dtype", "shape": (16, 16), "dtype": torch.float16, "dim": 1},
-        {"name": "wrong_output_shape", "shape": (16, 16), "dtype": torch.float32, "out_shape": (15,), "dim": 1},
-        {"name": "non_contiguous", "shape": (16, 16), "dtype": torch.float32, "dim": 1},
-        {"name": "wrong_rank", "shape": (16,), "dtype": torch.float32, "dim": 1},
-    ]
-
-
-def benchmark_cases():
-    return [
-        {"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/relu.py
+++ b/tests/cases/relu.py
@ -1,35 +0,0 @@
-from __future__ import annotations
-
-import torch
-
-
-def correctness_cases():
-    return [
-        {"name": "relu_1k_fp32", "shape": (1024,), "dtype": torch.float32, "negative_slope": 0.0, "atol": 0, "rtol": 0},
-        {"name": "relu_1k_fp16", "shape": (1024,), "dtype": torch.float16, "negative_slope": 0.0, "atol": 0, "rtol": 0},
-        {"name": "leaky_relu_2d_fp32", "shape": (32, 64), "dtype": torch.float32, "negative_slope": 0.01, "atol": 1e-6, "rtol": 1e-6},
-        {"name": "leaky_relu_2d_fp16", "shape": (32, 64), "dtype": torch.float16, "negative_slope": 0.01, "atol": 1e-3, "rtol": 1e-3},
-        {"name": "non_contiguous_fp32", "shape": (16, 32), "dtype": torch.float32, "negative_slope": 0.01, "atol": 1e-6, "rtol": 1e-6},
-        {"name": "broadcast_input_fp32", "src_shape": (64,), "out_shape": (32, 64), "dtype": torch.float32, "negative_slope": 0.01, "atol": 1e-6, "rtol": 1e-6},
-    ]
-
-
-def api_error_cases():
-    return [
-        {"name": "shape_mismatch", "shape": (16,), "out_shape": (8,), "dtype": torch.float32},
-        {"name": "dtype_mismatch", "shape": (16,), "dtype": torch.float32, "out_dtype": torch.float16},
-        {"name": "cpu_tensor", "shape": (16,), "dtype": torch.float32},
-        {"name": "unsupported_dtype", "shape": (16,), "dtype": torch.float64},
-        {"name": "broadcasted_output", "shape": (8, 16), "base_shape": (1, 16), "dtype": torch.float32},
-    ]
-
-
-def benchmark_cases():
-    return [
-        {"name": "contiguous_1001k_fp16", "shape": (1001 * 1024,), "dtype": torch.float16, "negative_slope": 0.0},
-        {"name": "contiguous_4093k_fp16", "shape": (4093 * 1024,), "dtype": torch.float16, "negative_slope": 0.0},
-        {"name": "contiguous_65521k_fp16", "shape": (65521 * 1024,), "dtype": torch.float16, "negative_slope": 0.0},
-        {"name": "contiguous_1001k_fp32", "shape": (1001 * 1024,), "dtype": torch.float32, "negative_slope": 0.01},
-        {"name": "contiguous_4093k_fp32", "shape": (4093 * 1024,), "dtype": torch.float32, "negative_slope": 0.01},
-        {"name": "contiguous_65521k_fp32", "shape": (65521 * 1024,), "dtype": torch.float32, "negative_slope": 0.01},
-    ]
--- a/tests/cases/softmax.py
+++ b/tests/cases/softmax.py
@ -1,33 +0,0 @@
-from __future__ import annotations
-
-import torch
-
-
-def correctness_cases():
-    return [
-        {"name": "rowwise_1x128", "shape": (1, 128), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-        {"name": "rowwise_16x128", "shape": (16, 128), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-        {"name": "rowwise_16x256", "shape": (16, 256), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-        {"name": "rowwise_32x128", "shape": (32, 128), "dtype": torch.float32, "atol": 1e-5, "rtol": 1e-5},
-    ]
-
-
-def api_error_cases():
-    return [
-        {"name": "wrong_dim", "shape": (16, 16), "dtype": torch.float32, "dim": 0},
-        {"name": "wrong_dtype", "shape": (16, 16), "dtype": torch.float16, "dim": 1},
-        {"name": "wrong_output_shape", "shape": (16, 16), "dtype": torch.float32, "out_shape": (16, 15), "dim": 1},
-        {"name": "non_contiguous", "shape": (16, 16), "dtype": torch.float32, "dim": 1},
-        {"name": "wrong_rank", "shape": (16,), "dtype": torch.float32, "dim": 1},
-    ]
-
-
-def benchmark_cases():
-    return [
-        {"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
@ -1,31 +0,0 @@
-from __future__ import annotations
-
-import torch
-
-
-def correctness_cases():
-    return [
-        {"name": "contiguous_1k_fp32", "shape": (1024,), "dtype": torch.float32, "atol": 1e-6, "rtol": 1e-6},
-        {"name": "contiguous_1k_fp16", "shape": (1024,), "dtype": torch.float16, "atol": 1e-3, "rtol": 1e-3},
-        {"name": "contiguous_1536_fp32", "shape": (1536,), "dtype": torch.float32, "atol": 1e-6, "rtol": 1e-6},
-        {"name": "contiguous_1536_fp16", "shape": (1536,), "dtype": torch.float16, "atol": 1e-3, "rtol": 1e-3},
-    ]
-
-
-def api_error_cases():
-    return [
-        {"name": "shape_mismatch", "shape": (16,), "other_shape": (8,), "dtype": torch.float32},
-        {"name": "dtype_mismatch", "shape": (16,), "dtype": torch.float32, "other_dtype": torch.float16},
-        {"name": "non_contiguous", "shape": (4, 4), "dtype": torch.float32},
-    ]
-
-
-def benchmark_cases():
-    return [
-        {"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/conftest.py
+++ b/tests/conftest.py
@ -1,24 +0,0 @@
-from __future__ import annotations
-
-import os
-import sys
-from pathlib import Path
-
-import pytest
-
-ROOT = Path(__file__).resolve().parents[1]
-PYTHON_DIR = ROOT / "python"
-if str(PYTHON_DIR) not in sys.path:
-    sys.path.insert(0, str(PYTHON_DIR))
-if str(ROOT) not in sys.path:
-    sys.path.insert(0, str(ROOT))
-
-
-def pytest_addoption(parser):
-    parser.addoption("--backend", action="store", default=os.environ.get("CAMP_TEST_BACKEND", "nvidia"))
-
-
-@pytest.fixture
-def backend(request) -> str:
-    return request.config.getoption("--backend")
-
--- a/Show More
+++ b/Show More