feat: scaffold operator runtime backends

Co-authored-by: wawahejun <hejunlbbc@gmail.com>

CMakeLists.txt

@@ -0,0 +1,25 @@
+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(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+
+include(cmake/GenerateOperators.cmake)
+
+if(CAMP_ENABLE_NVIDIA)
+  if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+    set(CMAKE_CUDA_ARCHITECTURES native)
+  endif()
+  enable_language(CUDA)
+  include(cmake/cuda_helpers.cmake)
+endif()
+
+add_subdirectory(ops)

cmake/FindMetaX.cmake

@@ -0,0 +1,6 @@
+# Placeholder for a future MetaX toolchain integration.
+#
+# A production implementation should detect MACA_PATH, htcc/mxcc, runtime
+# libraries, include directories, and backend-specific source suffixes.
+set(MetaX_FOUND OFF)
+

cmake/GenerateOperators.cmake

@@ -0,0 +1,22 @@
+set(CAMP_GENERATED_DIR "${CMAKE_BINARY_DIR}/generated")
+file(MAKE_DIRECTORY "${CAMP_GENERATED_DIR}")
+
+find_package(Python3 REQUIRED COMPONENTS Interpreter)
+
+set(CAMP_GENERATED_OPERATORS_CMAKE "${CAMP_GENERATED_DIR}/operators.cmake")
+set(CAMP_GENERATED_REGISTRY_PY "${CAMP_GENERATED_DIR}/operator_registry.py")
+set(CAMP_GENERATED_TEST_MANIFEST "${CAMP_GENERATED_DIR}/operator_test_manifest.json")
+
+execute_process(
+  COMMAND "${Python3_EXECUTABLE}"
+          "${CMAKE_CURRENT_SOURCE_DIR}/tools/generate_operator_artifacts.py"
+          --ops-root "${CMAKE_CURRENT_SOURCE_DIR}/ops"
+          --out-dir "${CAMP_GENERATED_DIR}"
+  WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}"
+  RESULT_VARIABLE CAMP_GENERATOR_RESULT
+)
+
+if(NOT CAMP_GENERATOR_RESULT EQUAL 0)
+  message(FATAL_ERROR "operator artifact generation failed")
+endif()
+

cmake/cuda_helpers.cmake

@@ -0,0 +1,10 @@
+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)
+

include/operator_runtime/api.h

@@ -0,0 +1,48 @@
+#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;
+
+#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);
+
+#ifdef __cplusplus
+}
+#endif
+

include/operator_runtime/cuda_helpers.h

@@ -0,0 +1,30 @@
+#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
+

include/operator_runtime/descriptor.h

@@ -0,0 +1,16 @@
+#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;
+};
+
+#endif
+

include/operator_runtime/elementwise.h

@@ -0,0 +1,24 @@
+#pragma once
+
+#include "operator_runtime/tensor_view.h"
+
+#ifdef __cplusplus
+
+namespace oprt {
+
+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);
+}
+
+} // namespace oprt
+
+#endif
+

include/operator_runtime/operation.h

@@ -0,0 +1,18 @@
+#pragma once
+
+#ifdef __cplusplus
+
+#include <string>
+
+namespace oprt {
+
+struct OperationSpec {
+    std::string name;
+    std::string backend;
+    std::string kind;
+};
+
+} // namespace oprt
+
+#endif
+

include/operator_runtime/tensor_checks.h

@@ -0,0 +1,40 @@
+#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
+

include/operator_runtime/tensor_view.h

@@ -0,0 +1,55 @@
+#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;
+}
+
+} // namespace oprt
+
+#endif
+

ops/CMakeLists.txt

@@ -0,0 +1,17 @@
+include("${CMAKE_BINARY_DIR}/generated/operators.cmake")
+
+add_library(camp_ops SHARED ${CAMP_OPERATOR_SOURCES})
+target_include_directories(camp_ops PUBLIC
+  "${CMAKE_CURRENT_SOURCE_DIR}/.."
+  "${CMAKE_CURRENT_SOURCE_DIR}/../include"
+)
+
+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()
+

ops/common/metax/README.md

@@ -0,0 +1,13 @@
+# MetaX Backend Stub
+
+The training runtime keeps MetaX interfaces aligned with NVIDIA interfaces, but
+does not build MetaX code by default.
+
+A production integration should add:
+
+- MACA SDK discovery (`MACA_PATH`)
+- htcc/mxcc compiler rules
+- `.maca` source handling
+- hcruntime/mcruntime link rules
+- backend-specific CI on MetaX hardware
+

ops/common/nvidia/elementwise.cuh

@@ -0,0 +1,45 @@
+#pragma once
+
+#include <cuda_runtime.h>
+#include <cuda_fp16.h>
+#include <stdint.h>
+
+namespace oprt::nvidia {
+
+template <typename T, typename UnaryOp>
+__global__ void unary_contiguous_kernel(T *out, const T *in, int64_t n, UnaryOp op) {
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t stride = int64_t(blockDim.x) * gridDim.x;
+    for (int64_t i = idx; i < n; i += stride) {
+        out[i] = op(in[i]);
+    }
+}
+
+template <typename T, typename BinaryOp>
+__global__ void binary_contiguous_kernel(T *out, const T *a, const T *b, int64_t n, BinaryOp op) {
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t stride = int64_t(blockDim.x) * gridDim.x;
+    for (int64_t i = idx; i < n; i += stride) {
+        out[i] = op(a[i], b[i]);
+    }
+}
+
+struct CopyOp {
+    template <typename T>
+    __device__ T operator()(T value) const {
+        return value;
+    }
+};
+
+struct AddOp {
+    __device__ float operator()(float a, float b) const {
+        return a + b;
+    }
+
+    __device__ half operator()(half a, half b) const {
+        return __hadd(a, b);
+    }
+};
+
+} // namespace oprt::nvidia
+

ops/common/nvidia/indexing.cuh

@@ -0,0 +1,12 @@
+#pragma once
+
+#include <stdint.h>
+
+namespace oprt::nvidia {
+
+__device__ inline int64_t contiguous_offset(int64_t linear) {
+    return linear;
+}
+
+} // namespace oprt::nvidia
+

ops/common/status.cc

@@ -0,0 +1,21 @@
+#include "operator_runtime/api.h"
+
+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";
+    }
+}
+

ops/copy/metax/README.md

@@ -0,0 +1,5 @@
+# Copy MetaX Stub
+
+The copy ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and htcc/mxcc
+rules when MetaX hardware is available.
+

ops/copy/metax/copy_metax.h

@@ -0,0 +1,6 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
+

ops/copy/nvidia/copy_cuda.cu

@@ -0,0 +1,112 @@
+#include "ops/copy/nvidia/copy_cuda.h"
+
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/elementwise.h"
+#include "operator_runtime/tensor_checks.h"
+#include "operator_runtime/cuda_helpers.h"
+#include "ops/common/nvidia/elementwise.cuh"
+
+#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;
+    bool fast_path = false;
+
+    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);
+    oprt::nvidia::unary_contiguous_kernel<T><<<blocks, threads, 0, oprt::as_cuda_stream(stream)>>>(
+        static_cast<T *>(dst), static_cast<const T *>(src), desc->elements, oprt::nvidia::CopyOp{});
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+} // namespace
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia(
+    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->fast_path = true;
+    typed->workspace_size = 0;
+    *desc = typed;
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia(
+    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_nvidia(
+    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_nvidia(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}

ops/copy/nvidia/copy_cuda.h

@@ -0,0 +1,32 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *dst,
+    const oprt_tensor_view_t *src);
+
+OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia(
+    oprt_operator_descriptor_t desc,
+    size_t *size);
+
+OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia(
+    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_nvidia(
+    oprt_operator_descriptor_t desc);
+
+#ifdef __cplusplus
+}
+#endif
+

ops/copy/operator.yaml

@@ -0,0 +1,46 @@
+name: copy
+kind: elementwise_unary
+python_module: operator_runtime.ops.copy
+torch_reference: torch.clone
+backends:
+  nvidia:
+    enabled_by_default: true
+    status: runnable
+    sources:
+      - ops/copy/nvidia/copy_cuda.cu
+    headers:
+      - ops/copy/nvidia/copy_cuda.h
+    symbols:
+      create: oprt_create_copy_descriptor_nvidia
+      workspace: oprt_get_copy_workspace_size_nvidia
+      execute: oprt_execute_copy_nvidia
+      destroy: oprt_destroy_copy_descriptor_nvidia
+    dtypes: [float16, float32]
+    supports:
+      broadcast: false
+      strided: false
+  tilelang:
+    enabled_by_default: true
+    status: runnable
+    module: ops.copy.tilelang.copy_tl
+    dtypes: [float16, float32]
+    supports:
+      broadcast: false
+      strided: false
+  metax:
+    enabled_by_default: false
+    status: stub
+    headers:
+      - ops/copy/metax/copy_metax.h
+tolerances:
+  float16: {atol: 0, rtol: 0}
+  float32: {atol: 0, rtol: 0}
+benchmark:
+  against_torch: true
+  default_cases:
+    - name: contiguous_1m
+performance_model:
+  bound: memory
+  bytes: "2 * numel * elem_bytes"
+  flops: "0"
+

ops/copy/tilelang/copy_tl.py

@@ -0,0 +1,79 @@
+from __future__ import annotations
+
+from functools import lru_cache
+from dataclasses import dataclass
+
+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_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(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(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_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)

ops/reduce_sum/metax/README.md

@@ -0,0 +1,5 @@
+# Reduce Sum MetaX Stub
+
+The reduce_sum ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and
+htcc/mxcc rules when MetaX hardware is available.
+

ops/reduce_sum/metax/reduce_sum_metax.h

@@ -0,0 +1,6 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
+

ops/reduce_sum/nvidia/reduce_sum_cuda.cu

@@ -0,0 +1,126 @@
+#include "ops/reduce_sum/nvidia/reduce_sum_cuda.h"
+
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/tensor_checks.h"
+#include "operator_runtime/cuda_helpers.h"
+
+#include <cuda_runtime.h>
+#include <float.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) {
+    extern __shared__ float smem[];
+    int row = blockIdx.x;
+    float sum = 0.0f;
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        sum += in[int64_t(row) * cols + col];
+    }
+    smem[threadIdx.x] = sum;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
+        if (threadIdx.x < stride) {
+            smem[threadIdx.x] += smem[threadIdx.x + stride];
+        }
+        __syncthreads();
+    }
+    if (threadIdx.x == 0) {
+        out[row] = smem[0];
+    }
+}
+
+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_nvidia(
+    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_nvidia(
+    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_nvidia(
+    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);
+    constexpr int threads = 256;
+    reduce_sum_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), oprt::as_cuda_stream(stream)>>>(
+        static_cast<float *>(out), static_cast<const float *>(in), typed->rows, typed->cols);
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor_nvidia(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}
+

ops/reduce_sum/nvidia/reduce_sum_cuda.h

@@ -0,0 +1,33 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *out,
+    const oprt_tensor_view_t *in,
+    int64_t axis);
+
+OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia(
+    oprt_operator_descriptor_t desc,
+    size_t *size);
+
+OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia(
+    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_nvidia(
+    oprt_operator_descriptor_t desc);
+
+#ifdef __cplusplus
+}
+#endif
+

ops/reduce_sum/operator.yaml

@@ -0,0 +1,45 @@
+name: reduce_sum
+kind: reduction
+python_module: operator_runtime.ops.reduce_sum
+torch_reference: torch.sum
+backends:
+  nvidia:
+    enabled_by_default: true
+    status: runnable
+    sources:
+      - ops/reduce_sum/nvidia/reduce_sum_cuda.cu
+    headers:
+      - ops/reduce_sum/nvidia/reduce_sum_cuda.h
+    symbols:
+      create: oprt_create_reduce_sum_descriptor_nvidia
+      workspace: oprt_get_reduce_sum_workspace_size_nvidia
+      execute: oprt_execute_reduce_sum_nvidia
+      destroy: oprt_destroy_reduce_sum_descriptor_nvidia
+    dtypes: [float32]
+    supports:
+      broadcast: false
+      strided: false
+  tilelang:
+    enabled_by_default: true
+    status: runnable
+    module: ops.reduce_sum.tilelang.reduce_sum_tl
+    dtypes: [float32]
+    supports:
+      broadcast: false
+      strided: false
+  metax:
+    enabled_by_default: false
+    status: stub
+    headers:
+      - ops/reduce_sum/metax/reduce_sum_metax.h
+tolerances:
+  float32: {atol: 1e-5, rtol: 1e-5}
+benchmark:
+  against_torch: true
+  default_cases:
+    - name: rowwise_1024x1024
+performance_model:
+  bound: memory
+  bytes: "(numel + rows) * elem_bytes"
+  flops: "numel"
+

ops/reduce_sum/tilelang/reduce_sum_tl.py

@@ -0,0 +1,88 @@
+from __future__ import annotations
+
+from dataclasses import dataclass
+from functools import lru_cache
+
+import tilelang
+import tilelang.language as T
+import torch
+
+
+@tilelang.jit(
+    pass_configs={
+        tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True,
+        tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: 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)
+
+    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
+        src_local = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
+        out_local = T.alloc_fragment((BLOCK_N,), dtype)
+        T.clear(out_local)
+
+        for m_blk in T.Serial(M // BLOCK_M):
+            T.copy(src[pid_n * BLOCK_N, m_blk * BLOCK_M], src_local)
+            T.reduce_sum(src_local, out_local, dim=1, clear=False)
+
+        T.copy(out_local, out[pid_n * BLOCK_N])
+
+    return out
+
+
+@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(self) -> None:
+        assert self.kernel is not 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_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)

ops/softmax/metax/README.md

@@ -0,0 +1,5 @@
+# Softmax MetaX Stub
+
+The softmax ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and htcc/mxcc
+rules when MetaX hardware is available.
+

ops/softmax/metax/softmax_metax.h

@@ -0,0 +1,6 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
+

ops/softmax/nvidia/softmax_cuda.cu

@@ -0,0 +1,147 @@
+#include "ops/softmax/nvidia/softmax_cuda.h"
+
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/tensor_checks.h"
+#include "operator_runtime/cuda_helpers.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) {
+    extern __shared__ float smem[];
+    int row = blockIdx.x;
+    float local_max = -FLT_MAX;
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        float value = in[int64_t(row) * cols + col];
+        local_max = fmaxf(local_max, value);
+    }
+    smem[threadIdx.x] = local_max;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
+        if (threadIdx.x < stride) {
+            smem[threadIdx.x] = fmaxf(smem[threadIdx.x], smem[threadIdx.x + stride]);
+        }
+        __syncthreads();
+    }
+    float row_max = smem[0];
+
+    float local_sum = 0.0f;
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        float value = expf(in[int64_t(row) * cols + col] - row_max);
+        out[int64_t(row) * cols + col] = value;
+        local_sum += value;
+    }
+    smem[threadIdx.x] = local_sum;
+    __syncthreads();
+
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
+        if (threadIdx.x < stride) {
+            smem[threadIdx.x] += smem[threadIdx.x + stride];
+        }
+        __syncthreads();
+    }
+    float row_sum = smem[0];
+
+    for (int64_t col = threadIdx.x; col < cols; col += blockDim.x) {
+        out[int64_t(row) * cols + col] /= row_sum;
+    }
+}
+
+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_nvidia(
+    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_nvidia(
+    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_nvidia(
+    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);
+    constexpr int threads = 256;
+    softmax_rowwise_kernel<<<typed->rows, threads, threads * sizeof(float), oprt::as_cuda_stream(stream)>>>(
+        static_cast<float *>(out), static_cast<const float *>(in), typed->rows, typed->cols);
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_nvidia(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}
+

ops/softmax/nvidia/softmax_cuda.h

@@ -0,0 +1,33 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *out,
+    const oprt_tensor_view_t *in,
+    int64_t axis);
+
+OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia(
+    oprt_operator_descriptor_t desc,
+    size_t *size);
+
+OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia(
+    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_nvidia(
+    oprt_operator_descriptor_t desc);
+
+#ifdef __cplusplus
+}
+#endif
+

ops/softmax/operator.yaml

@@ -0,0 +1,45 @@
+name: softmax
+kind: normalization
+python_module: operator_runtime.ops.softmax
+torch_reference: torch.softmax
+backends:
+  nvidia:
+    enabled_by_default: true
+    status: runnable
+    sources:
+      - ops/softmax/nvidia/softmax_cuda.cu
+    headers:
+      - ops/softmax/nvidia/softmax_cuda.h
+    symbols:
+      create: oprt_create_softmax_descriptor_nvidia
+      workspace: oprt_get_softmax_workspace_size_nvidia
+      execute: oprt_execute_softmax_nvidia
+      destroy: oprt_destroy_softmax_descriptor_nvidia
+    dtypes: [float32]
+    supports:
+      broadcast: false
+      strided: false
+  tilelang:
+    enabled_by_default: true
+    status: runnable
+    module: ops.softmax.tilelang.softmax_tl
+    dtypes: [float32]
+    supports:
+      broadcast: false
+      strided: false
+  metax:
+    enabled_by_default: false
+    status: stub
+    headers:
+      - ops/softmax/metax/softmax_metax.h
+tolerances:
+  float32: {atol: 1e-5, rtol: 1e-5}
+benchmark:
+  against_torch: true
+  default_cases:
+    - name: rowwise_1024x1024
+performance_model:
+  bound: mixed
+  bytes: "5 * numel * elem_bytes"
+  flops: "4 * numel"
+

ops/softmax/tilelang/softmax_tl.py

@@ -0,0 +1,108 @@
+from __future__ import annotations
+
+from dataclasses import dataclass
+from functools import lru_cache
+
+import tilelang
+import tilelang.language as T
+import torch
+
+
+@tilelang.jit(
+    pass_configs={
+        tilelang.PassConfigKey.TL_DISABLE_WARP_SPECIALIZED: True,
+        tilelang.PassConfigKey.TL_DISABLE_TMA_LOWER: 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)
+
+    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
+        src_local = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
+        out_local = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
+        cur_exp = T.alloc_fragment((BLOCK_N, BLOCK_M), dtype)
+        cur_max = T.alloc_fragment((BLOCK_N,), dtype)
+        cur_sum = T.alloc_fragment((BLOCK_N,), dtype)
+        lse = T.alloc_fragment((BLOCK_N,), dtype)
+
+        T.fill(lse, -T.infinity(dtype))
+
+        for m_blk in T.Serial(M // BLOCK_M):
+            T.copy(src[pid_n * BLOCK_N, m_blk * BLOCK_M], src_local)
+            T.reduce_max(src_local, cur_max, dim=1, clear=True)
+
+            for i, j in T.Parallel(BLOCK_N, BLOCK_M):
+                cur_exp[i, j] = T.exp2(src_local[i, j] * log2_e - cur_max[i] * log2_e)
+
+            T.reduce_sum(cur_exp, cur_sum, dim=1, clear=True)
+
+            for i in T.Parallel(BLOCK_N):
+                lse[i] = cur_max[i] * log2_e + T.log2(
+                    T.exp2(lse[i] - cur_max[i] * log2_e) + cur_sum[i]
+                )
+
+        for m_blk in T.Serial(M // BLOCK_M):
+            T.copy(src[pid_n * BLOCK_N, m_blk * BLOCK_M], src_local)
+            for i, j in T.Parallel(BLOCK_N, BLOCK_M):
+                out_local[i, j] = T.exp2(src_local[i, j] * log2_e - lse[i])
+            T.copy(out_local, out[pid_n * BLOCK_N, m_blk * BLOCK_M])
+
+    return out
+
+
+@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(self) -> None:
+        assert self.kernel is not 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_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)

ops/vector_add/metax/README.md

@@ -0,0 +1,5 @@
+# Vector Add MetaX Stub
+
+The vector_add ABI mirrors the NVIDIA lifecycle. Add `.maca` sources and
+htcc/mxcc rules when MetaX hardware is available.
+

ops/vector_add/metax/vector_add_metax.h

@@ -0,0 +1,6 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+// MetaX stub: keep the ABI contract visible while the real toolchain is absent.
+

ops/vector_add/nvidia/vector_add_cuda.cu

@@ -0,0 +1,118 @@
+#include "ops/vector_add/nvidia/vector_add_cuda.h"
+
+#include "operator_runtime/descriptor.h"
+#include "operator_runtime/elementwise.h"
+#include "operator_runtime/tensor_checks.h"
+#include "operator_runtime/cuda_helpers.h"
+#include "ops/common/nvidia/elementwise.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);
+    oprt::nvidia::binary_contiguous_kernel<T><<<blocks, threads, 0, oprt::as_cuda_stream(stream)>>>(
+        static_cast<T *>(out), static_cast<const T *>(a), static_cast<const T *>(b), desc->elements, oprt::nvidia::AddOp{});
+    OPRT_CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return OPRT_SUCCESS;
+}
+
+} // namespace
+
+extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia(
+    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_nvidia(
+    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_nvidia(
+    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_nvidia(
+    oprt_operator_descriptor_t desc) {
+    delete desc;
+    return OPRT_SUCCESS;
+}

ops/vector_add/nvidia/vector_add_cuda.h

@@ -0,0 +1,34 @@
+#pragma once
+
+#include "operator_runtime/api.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia(
+    oprt_operator_descriptor_t *desc,
+    const oprt_tensor_view_t *out,
+    const oprt_tensor_view_t *a,
+    const oprt_tensor_view_t *b);
+
+OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_nvidia(
+    oprt_operator_descriptor_t desc,
+    size_t *size);
+
+OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia(
+    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_nvidia(
+    oprt_operator_descriptor_t desc);
+
+#ifdef __cplusplus
+}
+#endif
+

ops/vector_add/operator.yaml

@@ -0,0 +1,46 @@
+name: vector_add
+kind: elementwise_binary
+python_module: operator_runtime.ops.vector_add
+torch_reference: torch.add
+backends:
+  nvidia:
+    enabled_by_default: true
+    status: runnable
+    sources:
+      - ops/vector_add/nvidia/vector_add_cuda.cu
+    headers:
+      - ops/vector_add/nvidia/vector_add_cuda.h
+    symbols:
+      create: oprt_create_vector_add_descriptor_nvidia
+      workspace: oprt_get_vector_add_workspace_size_nvidia
+      execute: oprt_execute_vector_add_nvidia
+      destroy: oprt_destroy_vector_add_descriptor_nvidia
+    dtypes: [float16, float32]
+    supports:
+      broadcast: false
+      strided: false
+  tilelang:
+    enabled_by_default: true
+    status: runnable
+    module: ops.vector_add.tilelang.vector_add_tl
+    dtypes: [float16, float32]
+    supports:
+      broadcast: false
+      strided: false
+  metax:
+    enabled_by_default: false
+    status: stub
+    headers:
+      - ops/vector_add/metax/vector_add_metax.h
+tolerances:
+  float16: {atol: 1e-3, rtol: 1e-3}
+  float32: {atol: 1e-6, rtol: 1e-6}
+benchmark:
+  against_torch: true
+  default_cases:
+    - name: contiguous_1m
+performance_model:
+  bound: memory
+  bytes: "3 * numel * elem_bytes"
+  flops: "numel"
+

ops/vector_add/tilelang/vector_add_tl.py

@@ -0,0 +1,80 @@
+from __future__ import annotations
+
+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 _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)
+
+    with T.Kernel(N // BLOCK_N, threads=256) as pid_n:
+        base = pid_n * BLOCK_N
+        for i in T.Parallel(BLOCK_N):
+            out[base + i] = a[base + i] + b[base + i]
+
+    return out
+
+
+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(self) -> None:
+        self.out.copy_(self.kernel(self.a, self.b))
+
+    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)

python/operator_runtime/__init__.py

@@ -0,0 +1,23 @@
+from .backend import Backend, normalize_backend
+from .ops.copy import copy, copy_, prepare_copy
+from .ops.reduce_sum import prepare_reduce_sum, reduce_sum, reduce_sum_
+from .ops.softmax import prepare_softmax, softmax, softmax_
+from .ops.vector_add import prepare_vector_add, vector_add, vector_add_
+
+__all__ = [
+    "Backend",
+    "normalize_backend",
+    "copy",
+    "copy_",
+    "prepare_copy",
+    "vector_add",
+    "vector_add_",
+    "prepare_vector_add",
+    "reduce_sum",
+    "reduce_sum_",
+    "prepare_reduce_sum",
+    "softmax",
+    "softmax_",
+    "prepare_softmax",
+]
+

python/operator_runtime/backend.py

@@ -0,0 +1,19 @@
+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
+

python/operator_runtime/benchmark.py

@@ -0,0 +1,20 @@
+from __future__ import annotations
+
+from collections.abc import Callable
+
+import torch
+
+
+def cuda_time_ms(fn: Callable[[], None], *, warmup: int = 10, iterations: int = 100) -> float:
+    for _ in range(warmup):
+        fn()
+    torch.cuda.synchronize()
+    start = torch.cuda.Event(enable_timing=True)
+    end = torch.cuda.Event(enable_timing=True)
+    start.record()
+    for _ in range(iterations):
+        fn()
+    end.record()
+    torch.cuda.synchronize()
+    return start.elapsed_time(end) / iterations
+

python/operator_runtime/ctypes_bindings.py

@@ -0,0 +1,126 @@
+from __future__ import annotations
+
+import ctypes
+from dataclasses import dataclass
+from typing import Callable
+
+from .loader import load_library
+from .tensor_view import TensorView
+
+Status = ctypes.c_int
+Descriptor = ctypes.c_void_p
+
+
+class OperatorRuntimeError(RuntimeError):
+    pass
+
+
+def check_status(status: int) -> None:
+    if status == 0:
+        return
+    lib = load_library()
+    lib.oprt_status_string.argtypes = [ctypes.c_int]
+    lib.oprt_status_string.restype = ctypes.c_char_p
+    message = lib.oprt_status_string(status).decode("utf-8")
+    raise OperatorRuntimeError(message)
+
+
+@dataclass(frozen=True)
+class CFunctions:
+    create: Callable
+    workspace: Callable
+    execute: Callable
+    destroy: Callable
+
+
+def bind_unary(name: str) -> CFunctions:
+    lib = load_library()
+    create = getattr(lib, f"oprt_create_{name}_descriptor_nvidia")
+    create.argtypes = [ctypes.POINTER(Descriptor), ctypes.POINTER(TensorView), ctypes.POINTER(TensorView)]
+    create.restype = Status
+
+    workspace = getattr(lib, f"oprt_get_{name}_workspace_size_nvidia")
+    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
+    workspace.restype = Status
+
+    execute = getattr(lib, f"oprt_execute_{name}_nvidia")
+    execute.argtypes = [
+        Descriptor,
+        ctypes.c_void_p,
+        ctypes.c_size_t,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+    ]
+    execute.restype = Status
+
+    destroy = getattr(lib, f"oprt_destroy_{name}_descriptor_nvidia")
+    destroy.argtypes = [Descriptor]
+    destroy.restype = Status
+    return CFunctions(create, workspace, execute, destroy)
+
+
+def bind_binary(name: str) -> CFunctions:
+    lib = load_library()
+    create = getattr(lib, f"oprt_create_{name}_descriptor_nvidia")
+    create.argtypes = [
+        ctypes.POINTER(Descriptor),
+        ctypes.POINTER(TensorView),
+        ctypes.POINTER(TensorView),
+        ctypes.POINTER(TensorView),
+    ]
+    create.restype = Status
+
+    workspace = getattr(lib, f"oprt_get_{name}_workspace_size_nvidia")
+    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
+    workspace.restype = Status
+
+    execute = getattr(lib, f"oprt_execute_{name}_nvidia")
+    execute.argtypes = [
+        Descriptor,
+        ctypes.c_void_p,
+        ctypes.c_size_t,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+    ]
+    execute.restype = Status
+
+    destroy = getattr(lib, f"oprt_destroy_{name}_descriptor_nvidia")
+    destroy.argtypes = [Descriptor]
+    destroy.restype = Status
+    return CFunctions(create, workspace, execute, destroy)
+
+
+def bind_reduce_like(name: str) -> CFunctions:
+    lib = load_library()
+    create = getattr(lib, f"oprt_create_{name}_descriptor_nvidia")
+    create.argtypes = [
+        ctypes.POINTER(Descriptor),
+        ctypes.POINTER(TensorView),
+        ctypes.POINTER(TensorView),
+        ctypes.c_int64,
+    ]
+    create.restype = Status
+
+    workspace = getattr(lib, f"oprt_get_{name}_workspace_size_nvidia")
+    workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)]
+    workspace.restype = Status
+
+    execute = getattr(lib, f"oprt_execute_{name}_nvidia")
+    execute.argtypes = [
+        Descriptor,
+        ctypes.c_void_p,
+        ctypes.c_size_t,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+        ctypes.c_void_p,
+    ]
+    execute.restype = Status
+
+    destroy = getattr(lib, f"oprt_destroy_{name}_descriptor_nvidia")
+    destroy.argtypes = [Descriptor]
+    destroy.restype = Status
+    return CFunctions(create, workspace, execute, destroy)
+

python/operator_runtime/generated/operator_registry.py

@@ -0,0 +1,10 @@
+# Placeholder overwritten by tools/generate_operator_artifacts.py during CMake configure.
+OPERATORS = []
+
+
+def get_operator(name: str):
+    for op in OPERATORS:
+        if op["name"] == name:
+            return op
+    raise KeyError(name)
+

python/operator_runtime/loader.py

@@ -0,0 +1,32 @@
+from __future__ import annotations
+
+import ctypes
+import os
+from functools import lru_cache
+from pathlib import Path
+
+
+def _candidate_library_paths() -> list[Path]:
+    candidates: list[Path] = []
+    build_dir = os.environ.get("CAMP_BUILD_DIR")
+    if build_dir:
+        root = Path(build_dir)
+        candidates.extend([root / "libcamp_ops.so", root / "ops" / "libcamp_ops.so"])
+    repo_root = Path(__file__).resolve().parents[2]
+    candidates.extend(
+        [
+            repo_root / "build" / "libcamp_ops.so",
+            repo_root / "build" / "ops" / "libcamp_ops.so",
+        ]
+    )
+    return candidates
+
+
+@lru_cache(maxsize=1)
+def load_library() -> ctypes.CDLL:
+    for path in _candidate_library_paths():
+        if path.exists():
+            return ctypes.CDLL(str(path))
+    searched = ", ".join(str(path) for path in _candidate_library_paths())
+    raise FileNotFoundError(f"libcamp_ops.so not found; searched: {searched}")
+

python/operator_runtime/ops/copy.py

@@ -0,0 +1,55 @@
+from __future__ import annotations
+
+import ctypes
+
+import torch
+
+from operator_runtime.backend import Backend, normalize_backend
+from operator_runtime.ctypes_bindings import Descriptor, bind_unary, check_status
+from operator_runtime.prepared import PreparedOp
+from operator_runtime.tensor_view import tensor_view
+
+
+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 is not Backend.NVIDIA:
+        raise NotImplementedError(f"backend {backend.value} is not runnable")
+
+    funcs = bind_unary("copy")
+    desc = Descriptor()
+    out_view = tensor_view(out)
+    src_view = tensor_view(src)
+    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(src_view)))
+    workspace_size = ctypes.c_size_t()
+    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
+    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
+    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(src.data_ptr()))
+    return PreparedOp(funcs, desc, workspace, args, out)
+
+
+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)
+

python/operator_runtime/ops/reduce_sum.py

@@ -0,0 +1,64 @@
+from __future__ import annotations
+
+import ctypes
+
+import torch
+
+from operator_runtime.backend import Backend, normalize_backend
+from operator_runtime.ctypes_bindings import Descriptor, bind_reduce_like, check_status
+from operator_runtime.prepared import PreparedOp
+from operator_runtime.tensor_view import tensor_view
+
+
+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 is not Backend.NVIDIA:
+        raise NotImplementedError(f"backend {backend.value} is not runnable")
+
+    funcs = bind_reduce_like("reduce_sum")
+    desc = Descriptor()
+    out_view = tensor_view(out)
+    src_view = tensor_view(src)
+    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(src_view), ctypes.c_int64(dim)))
+    workspace_size = ctypes.c_size_t()
+    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
+    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
+    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(src.data_ptr()))
+    return PreparedOp(funcs, desc, workspace, args, out)
+
+
+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)
+

python/operator_runtime/ops/softmax.py

@@ -0,0 +1,62 @@
+from __future__ import annotations
+
+import ctypes
+
+import torch
+
+from operator_runtime.backend import Backend, normalize_backend
+from operator_runtime.ctypes_bindings import Descriptor, bind_reduce_like, check_status
+from operator_runtime.prepared import PreparedOp
+from operator_runtime.tensor_view import tensor_view
+
+
+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 is not Backend.NVIDIA:
+        raise NotImplementedError(f"backend {backend.value} is not runnable")
+
+    funcs = bind_reduce_like("softmax")
+    desc = Descriptor()
+    out_view = tensor_view(out)
+    src_view = tensor_view(src)
+    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(src_view), ctypes.c_int64(dim)))
+    workspace_size = ctypes.c_size_t()
+    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
+    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
+    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(src.data_ptr()))
+    return PreparedOp(funcs, desc, workspace, args, out)
+
+
+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)
+

python/operator_runtime/ops/vector_add.py

@@ -0,0 +1,61 @@
+from __future__ import annotations
+
+import ctypes
+
+import torch
+
+from operator_runtime.backend import Backend, normalize_backend
+from operator_runtime.ctypes_bindings import Descriptor, bind_binary, check_status
+from operator_runtime.prepared import PreparedOp
+from operator_runtime.tensor_view import tensor_view
+
+
+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 is not Backend.NVIDIA:
+        raise NotImplementedError(f"backend {backend.value} is not runnable")
+
+    funcs = bind_binary("vector_add")
+    desc = Descriptor()
+    out_view = tensor_view(out)
+    a_view = tensor_view(a)
+    b_view = tensor_view(b)
+    check_status(funcs.create(ctypes.byref(desc), ctypes.byref(out_view), ctypes.byref(a_view), ctypes.byref(b_view)))
+    workspace_size = ctypes.c_size_t()
+    check_status(funcs.workspace(desc, ctypes.byref(workspace_size)))
+    workspace = torch.empty(workspace_size.value, dtype=torch.uint8, device=out.device) if workspace_size.value else None
+    args = (ctypes.c_void_p(out.data_ptr()), ctypes.c_void_p(a.data_ptr()), ctypes.c_void_p(b.data_ptr()))
+    return PreparedOp(funcs, desc, workspace, args, out)
+
+
+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)
+

python/operator_runtime/perf_model.py

@@ -0,0 +1,25 @@
+from __future__ import annotations
+
+import torch
+
+
+def elem_bytes(dtype: torch.dtype) -> int:
+    return torch.empty((), dtype=dtype).element_size()
+
+
+def estimate_copy(tensor: torch.Tensor) -> tuple[int, int]:
+    return 2 * tensor.numel() * elem_bytes(tensor.dtype), 0
+
+
+def estimate_vector_add(tensor: torch.Tensor) -> tuple[int, int]:
+    return 3 * tensor.numel() * elem_bytes(tensor.dtype), tensor.numel()
+
+
+def estimate_reduce_sum(tensor: torch.Tensor) -> tuple[int, int]:
+    rows = tensor.shape[0]
+    return (tensor.numel() + rows) * elem_bytes(tensor.dtype), tensor.numel()
+
+
+def estimate_softmax(tensor: torch.Tensor) -> tuple[int, int]:
+    return 5 * tensor.numel() * elem_bytes(tensor.dtype), 4 * tensor.numel()
+

python/operator_runtime/prepared.py

@@ -0,0 +1,49 @@
+from __future__ import annotations
+
+import ctypes
+from dataclasses import dataclass
+from typing import Any
+
+import torch
+
+from .ctypes_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:
+        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,
+            *self.runner_args,
+            ctypes.c_void_p(stream.cuda_stream),
+        )
+        check_status(status)
+
+    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
+

python/operator_runtime/profiler.py

@@ -0,0 +1,34 @@
+from __future__ import annotations
+
+from dataclasses import dataclass
+
+
+@dataclass
+class PerformanceResult:
+    operator: str
+    backend: str
+    shape: str
+    dtype: str
+    bytes: int
+    flops: int
+    runtime_ms: float
+    torch_ms: float | None = None
+
+    @property
+    def gbytes_per_sec(self) -> float:
+        if self.runtime_ms <= 0:
+            return 0.0
+        return self.bytes / (1024**3) / (self.runtime_ms / 1000.0)
+
+    @property
+    def gflops_per_sec(self) -> float:
+        if self.runtime_ms <= 0:
+            return 0.0
+        return self.flops / 1e9 / (self.runtime_ms / 1000.0)
+
+    @property
+    def speedup(self) -> float | None:
+        if self.torch_ms is None or self.runtime_ms <= 0:
+            return None
+        return self.torch_ms / self.runtime_ms
+

python/operator_runtime/tensor_view.py

@@ -0,0 +1,55 @@
+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)
+

python/operator_runtime/testing.py

@@ -0,0 +1,14 @@
+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)
+

tools/generate_operator_artifacts.py

@@ -0,0 +1,98 @@
+#!/usr/bin/env python3
+from __future__ import annotations
+
+import argparse
+import json
+from pathlib import Path
+from typing import Any
+
+try:
+    import yaml
+except ImportError as exc:  # pragma: no cover
+    raise SystemExit("PyYAML is required to generate operator artifacts") from exc
+
+
+def load_operator_manifests(ops_root: Path) -> list[dict[str, Any]]:
+    manifests: list[dict[str, Any]] = []
+    for path in sorted(ops_root.glob("*/operator.yaml")):
+        data = yaml.safe_load(path.read_text()) or {}
+        data["_path"] = str(path)
+        manifests.append(data)
+    return manifests
+
+
+def cmake_bool(name: str) -> str:
+    return f"${{{name}}}"
+
+
+def generate_operators_cmake(manifests: list[dict[str, Any]], ops_root: Path) -> str:
+    lines = [
+        "# Generated by tools/generate_operator_artifacts.py",
+        "set(CAMP_OPERATOR_SOURCES",
+        "  ${CMAKE_CURRENT_SOURCE_DIR}/common/status.cc",
+    ]
+    for manifest in manifests:
+        nvidia = manifest.get("backends", {}).get("nvidia", {})
+        if nvidia.get("status") != "runnable":
+            continue
+        for source in nvidia.get("sources", []):
+            relative = Path(source)
+            if relative.parts and relative.parts[0] == "ops":
+                relative = Path(*relative.parts[1:])
+            lines.append(f"  $<$<BOOL:{cmake_bool('CAMP_ENABLE_NVIDIA')}>:${{CMAKE_CURRENT_SOURCE_DIR}}/{relative.as_posix()}>")
+    lines.append(")")
+    lines.append("")
+    return "\n".join(lines)
+
+
+def python_literal(value: Any) -> str:
+    return repr(value)
+
+
+def generate_python_registry(manifests: list[dict[str, Any]]) -> str:
+    serializable = []
+    for manifest in manifests:
+        item = {key: value for key, value in manifest.items() if not key.startswith("_")}
+        serializable.append(item)
+    return "\n".join(
+        [
+            "# Generated by tools/generate_operator_artifacts.py",
+            "from __future__ import annotations",
+            "",
+            "OPERATORS = " + python_literal(serializable),
+            "",
+            "def get_operator(name: str):",
+            "    for op in OPERATORS:",
+            "        if op['name'] == name:",
+            "            return op",
+            "    raise KeyError(name)",
+            "",
+        ]
+    )
+
+
+def generate_test_manifest(manifests: list[dict[str, Any]]) -> str:
+    payload = [{key: value for key, value in manifest.items() if not key.startswith("_")} for manifest in manifests]
+    return json.dumps(payload, indent=2, sort_keys=True) + "\n"
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--ops-root", type=Path, required=True)
+    parser.add_argument("--out-dir", type=Path, required=True)
+    args = parser.parse_args()
+
+    manifests = load_operator_manifests(args.ops_root)
+    if not manifests:
+        raise SystemExit(f"no operator manifests found under {args.ops_root}")
+
+    args.out_dir.mkdir(parents=True, exist_ok=True)
+    (args.out_dir / "operators.cmake").write_text(generate_operators_cmake(manifests, args.ops_root))
+    (args.out_dir / "operator_registry.py").write_text(generate_python_registry(manifests))
+    (args.out_dir / "operator_test_manifest.json").write_text(generate_test_manifest(manifests))
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
+

tools/operator_schema.yaml

@@ -0,0 +1,25 @@
+required_fields:
+  - name
+  - kind
+  - python_module
+  - torch_reference
+  - backends
+  - tolerances
+  - benchmark
+backend_required_fields:
+  runnable:
+    - status
+  stub:
+    - status
+supported_backends:
+  - nvidia
+  - tilelang
+  - metax
+supported_status:
+  - runnable
+  - stub
+  - spec-only
+supported_dtypes:
+  - float16
+  - float32
+

tools/validate_operator_manifest.py

@@ -0,0 +1,112 @@
+#!/usr/bin/env python3
+from __future__ import annotations
+
+import argparse
+import importlib.util
+import sys
+from pathlib import Path
+from typing import Any
+
+try:
+    import yaml
+except ImportError as exc:  # pragma: no cover
+    raise SystemExit("PyYAML is required to validate operator manifests") from exc
+
+
+REQUIRED_FIELDS = {
+    "name",
+    "kind",
+    "python_module",
+    "torch_reference",
+    "backends",
+    "tolerances",
+    "benchmark",
+}
+
+NVIDIA_SYMBOLS = {"create", "workspace", "execute", "destroy"}
+
+
+def fail(message: str) -> None:
+    raise ValueError(message)
+
+
+def load_manifest(path: Path) -> dict[str, Any]:
+    data = yaml.safe_load(path.read_text()) or {}
+    missing = REQUIRED_FIELDS - set(data)
+    if missing:
+        fail(f"{path}: missing required fields {sorted(missing)}")
+    return data
+
+
+def check_path(root: Path, manifest_path: Path, relative: str) -> None:
+    path = root / relative
+    if not path.exists():
+        fail(f"{manifest_path}: referenced path does not exist: {relative}")
+
+
+def check_backend_paths(root: Path, manifest_path: Path, manifest: dict[str, Any]) -> None:
+    for backend_name, backend in manifest["backends"].items():
+        status = backend.get("status")
+        if status == "runnable":
+            for key in ("sources", "headers"):
+                for relative in backend.get(key, []):
+                    check_path(root, manifest_path, relative)
+        if backend_name == "nvidia" and status == "runnable":
+            symbols = backend.get("symbols", {})
+            missing = NVIDIA_SYMBOLS - set(symbols)
+            if missing:
+                fail(f"{manifest_path}: nvidia backend missing symbols {sorted(missing)}")
+
+
+def check_python_module(repo_root: Path, manifest_path: Path, manifest: dict[str, Any]) -> None:
+    module_name = manifest["python_module"]
+    module_path = repo_root / "python" / Path(*module_name.split(".")).with_suffix(".py")
+    if not module_path.exists():
+        fail(f"{manifest_path}: python module does not exist: {module_name}")
+    spec = importlib.util.spec_from_file_location(module_name, module_path)
+    if spec is None:
+        fail(f"{manifest_path}: cannot create import spec for {module_name}")
+
+
+def check_cases(tests_root: Path, manifest_path: Path, manifest: dict[str, Any]) -> None:
+    case_path = tests_root / "cases" / f"{manifest['name']}.py"
+    if not case_path.exists():
+        fail(f"{manifest_path}: missing test case file {case_path}")
+    text = case_path.read_text()
+    for token in ("correctness_cases", "api_error_cases", "benchmark_cases"):
+        if token not in text:
+            fail(f"{manifest_path}: {case_path} missing {token}")
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--ops-root", type=Path, required=True)
+    parser.add_argument("--tests-root", type=Path, required=True)
+    args = parser.parse_args()
+
+    repo_root = Path.cwd()
+    manifests = sorted(args.ops_root.glob("*/operator.yaml"))
+    if not manifests:
+        raise SystemExit(f"no operator manifests found under {args.ops_root}")
+
+    errors: list[str] = []
+    for manifest_path in manifests:
+        try:
+            manifest = load_manifest(manifest_path)
+            check_backend_paths(repo_root, manifest_path, manifest)
+            check_python_module(repo_root, manifest_path, manifest)
+            check_cases(args.tests_root, manifest_path, manifest)
+        except Exception as exc:  # noqa: BLE001
+            errors.append(str(exc))
+
+    if errors:
+        for error in errors:
+            print(error, file=sys.stderr)
+        return 1
+    print(f"validated {len(manifests)} operator manifests")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
+