From ce30e4b19539dfacf50ef6ecfecde05afd448ce6 Mon Sep 17 00:00:00 2001 From: yutianyu Date: Sat, 2 May 2026 00:42:20 +0800 Subject: [PATCH] feat: scaffold operator runtime backends Co-authored-by: wawahejun --- CMakeLists.txt | 25 +++ cmake/FindMetaX.cmake | 6 + cmake/GenerateOperators.cmake | 22 +++ cmake/cuda_helpers.cmake | 10 ++ include/operator_runtime/api.h | 48 ++++++ include/operator_runtime/cuda_helpers.h | 30 ++++ include/operator_runtime/descriptor.h | 16 ++ include/operator_runtime/elementwise.h | 24 +++ include/operator_runtime/operation.h | 18 +++ include/operator_runtime/tensor_checks.h | 40 +++++ include/operator_runtime/tensor_view.h | 55 +++++++ ops/CMakeLists.txt | 17 ++ ops/common/metax/README.md | 13 ++ ops/common/nvidia/elementwise.cuh | 45 ++++++ ops/common/nvidia/indexing.cuh | 12 ++ ops/common/status.cc | 21 +++ ops/copy/metax/README.md | 5 + ops/copy/metax/copy_metax.h | 6 + ops/copy/nvidia/copy_cuda.cu | 112 +++++++++++++ ops/copy/nvidia/copy_cuda.h | 32 ++++ ops/copy/operator.yaml | 46 ++++++ ops/copy/tilelang/copy_tl.py | 79 ++++++++++ ops/reduce_sum/metax/README.md | 5 + ops/reduce_sum/metax/reduce_sum_metax.h | 6 + ops/reduce_sum/nvidia/reduce_sum_cuda.cu | 126 +++++++++++++++ ops/reduce_sum/nvidia/reduce_sum_cuda.h | 33 ++++ ops/reduce_sum/operator.yaml | 45 ++++++ ops/reduce_sum/tilelang/reduce_sum_tl.py | 88 +++++++++++ ops/softmax/metax/README.md | 5 + ops/softmax/metax/softmax_metax.h | 6 + ops/softmax/nvidia/softmax_cuda.cu | 147 ++++++++++++++++++ ops/softmax/nvidia/softmax_cuda.h | 33 ++++ ops/softmax/operator.yaml | 45 ++++++ ops/softmax/tilelang/softmax_tl.py | 108 +++++++++++++ ops/vector_add/metax/README.md | 5 + ops/vector_add/metax/vector_add_metax.h | 6 + ops/vector_add/nvidia/vector_add_cuda.cu | 118 ++++++++++++++ ops/vector_add/nvidia/vector_add_cuda.h | 34 ++++ ops/vector_add/operator.yaml | 46 ++++++ ops/vector_add/tilelang/vector_add_tl.py | 80 ++++++++++ python/operator_runtime/__init__.py | 23 +++ python/operator_runtime/backend.py | 19 +++ python/operator_runtime/benchmark.py | 20 +++ python/operator_runtime/ctypes_bindings.py | 126 +++++++++++++++ .../generated/operator_registry.py | 10 ++ python/operator_runtime/loader.py | 32 ++++ python/operator_runtime/ops/copy.py | 55 +++++++ python/operator_runtime/ops/reduce_sum.py | 64 ++++++++ python/operator_runtime/ops/softmax.py | 62 ++++++++ python/operator_runtime/ops/vector_add.py | 61 ++++++++ python/operator_runtime/perf_model.py | 25 +++ python/operator_runtime/prepared.py | 49 ++++++ python/operator_runtime/profiler.py | 34 ++++ python/operator_runtime/tensor_view.py | 55 +++++++ python/operator_runtime/testing.py | 14 ++ tools/generate_operator_artifacts.py | 98 ++++++++++++ tools/operator_schema.yaml | 25 +++ tools/validate_operator_manifest.py | 112 +++++++++++++ 58 files changed, 2502 insertions(+) create mode 100644 CMakeLists.txt create mode 100644 cmake/FindMetaX.cmake create mode 100644 cmake/GenerateOperators.cmake create mode 100644 cmake/cuda_helpers.cmake create mode 100644 include/operator_runtime/api.h create mode 100644 include/operator_runtime/cuda_helpers.h create mode 100644 include/operator_runtime/descriptor.h create mode 100644 include/operator_runtime/elementwise.h create mode 100644 include/operator_runtime/operation.h create mode 100644 include/operator_runtime/tensor_checks.h create mode 100644 include/operator_runtime/tensor_view.h create mode 100644 ops/CMakeLists.txt create mode 100644 ops/common/metax/README.md create mode 100644 ops/common/nvidia/elementwise.cuh create mode 100644 ops/common/nvidia/indexing.cuh create mode 100644 ops/common/status.cc create mode 100644 ops/copy/metax/README.md create mode 100644 ops/copy/metax/copy_metax.h create mode 100644 ops/copy/nvidia/copy_cuda.cu create mode 100644 ops/copy/nvidia/copy_cuda.h create mode 100644 ops/copy/operator.yaml create mode 100644 ops/copy/tilelang/copy_tl.py create mode 100644 ops/reduce_sum/metax/README.md create mode 100644 ops/reduce_sum/metax/reduce_sum_metax.h create mode 100644 ops/reduce_sum/nvidia/reduce_sum_cuda.cu create mode 100644 ops/reduce_sum/nvidia/reduce_sum_cuda.h create mode 100644 ops/reduce_sum/operator.yaml create mode 100644 ops/reduce_sum/tilelang/reduce_sum_tl.py create mode 100644 ops/softmax/metax/README.md create mode 100644 ops/softmax/metax/softmax_metax.h create mode 100644 ops/softmax/nvidia/softmax_cuda.cu create mode 100644 ops/softmax/nvidia/softmax_cuda.h create mode 100644 ops/softmax/operator.yaml create mode 100644 ops/softmax/tilelang/softmax_tl.py create mode 100644 ops/vector_add/metax/README.md create mode 100644 ops/vector_add/metax/vector_add_metax.h create mode 100644 ops/vector_add/nvidia/vector_add_cuda.cu create mode 100644 ops/vector_add/nvidia/vector_add_cuda.h create mode 100644 ops/vector_add/operator.yaml create mode 100644 ops/vector_add/tilelang/vector_add_tl.py create mode 100644 python/operator_runtime/__init__.py create mode 100644 python/operator_runtime/backend.py create mode 100644 python/operator_runtime/benchmark.py create mode 100644 python/operator_runtime/ctypes_bindings.py create mode 100644 python/operator_runtime/generated/operator_registry.py create mode 100644 python/operator_runtime/loader.py create mode 100644 python/operator_runtime/ops/copy.py create mode 100644 python/operator_runtime/ops/reduce_sum.py create mode 100644 python/operator_runtime/ops/softmax.py create mode 100644 python/operator_runtime/ops/vector_add.py create mode 100644 python/operator_runtime/perf_model.py create mode 100644 python/operator_runtime/prepared.py create mode 100644 python/operator_runtime/profiler.py create mode 100644 python/operator_runtime/tensor_view.py create mode 100644 python/operator_runtime/testing.py create mode 100644 tools/generate_operator_artifacts.py create mode 100644 tools/operator_schema.yaml create mode 100644 tools/validate_operator_manifest.py diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..e730bd0 --- /dev/null +++ b/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) diff --git a/cmake/FindMetaX.cmake b/cmake/FindMetaX.cmake new file mode 100644 index 0000000..5b1709f --- /dev/null +++ b/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) + diff --git a/cmake/GenerateOperators.cmake b/cmake/GenerateOperators.cmake new file mode 100644 index 0000000..f4e6e04 --- /dev/null +++ b/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() + diff --git a/cmake/cuda_helpers.cmake b/cmake/cuda_helpers.cmake new file mode 100644 index 0000000..23a276b --- /dev/null +++ b/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) + diff --git a/include/operator_runtime/api.h b/include/operator_runtime/api.h new file mode 100644 index 0000000..0e5b65c --- /dev/null +++ b/include/operator_runtime/api.h @@ -0,0 +1,48 @@ +#pragma once + +#include +#include + +#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 + diff --git a/include/operator_runtime/cuda_helpers.h b/include/operator_runtime/cuda_helpers.h new file mode 100644 index 0000000..1fbac5d --- /dev/null +++ b/include/operator_runtime/cuda_helpers.h @@ -0,0 +1,30 @@ +#pragma once + +#include "operator_runtime/api.h" + +#ifdef __CUDACC__ +#include + +#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(stream); +} + +inline int blocks_for(int64_t n, int threads) { + int64_t blocks = (n + threads - 1) / threads; + return static_cast(blocks > 0 ? blocks : 1); +} + +} // namespace oprt + +#endif + diff --git a/include/operator_runtime/descriptor.h b/include/operator_runtime/descriptor.h new file mode 100644 index 0000000..5d98a06 --- /dev/null +++ b/include/operator_runtime/descriptor.h @@ -0,0 +1,16 @@ +#pragma once + +#include "operator_runtime/api.h" + +#ifdef __cplusplus + +#include + +struct oprt_operator_descriptor { + virtual ~oprt_operator_descriptor() = default; + virtual const char *op_name() const = 0; + size_t workspace_size = 0; +}; + +#endif + diff --git a/include/operator_runtime/elementwise.h b/include/operator_runtime/elementwise.h new file mode 100644 index 0000000..f7d42ed --- /dev/null +++ b/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 + diff --git a/include/operator_runtime/operation.h b/include/operator_runtime/operation.h new file mode 100644 index 0000000..88a61e2 --- /dev/null +++ b/include/operator_runtime/operation.h @@ -0,0 +1,18 @@ +#pragma once + +#ifdef __cplusplus + +#include + +namespace oprt { + +struct OperationSpec { + std::string name; + std::string backend; + std::string kind; +}; + +} // namespace oprt + +#endif + diff --git a/include/operator_runtime/tensor_checks.h b/include/operator_runtime/tensor_checks.h new file mode 100644 index 0000000..47063df --- /dev/null +++ b/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 + diff --git a/include/operator_runtime/tensor_view.h b/include/operator_runtime/tensor_view.h new file mode 100644 index 0000000..367f908 --- /dev/null +++ b/include/operator_runtime/tensor_view.h @@ -0,0 +1,55 @@ +#pragma once + +#include "operator_runtime/api.h" + +#ifdef __cplusplus + +#include +#include + +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 + diff --git a/ops/CMakeLists.txt b/ops/CMakeLists.txt new file mode 100644 index 0000000..632b8ff --- /dev/null +++ b/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 + $<$:CAMP_ENABLE_NVIDIA=1> + $<$:CAMP_ENABLE_METAX=1> +) + +if(CAMP_ENABLE_NVIDIA) + set_target_properties(camp_ops PROPERTIES CUDA_SEPARABLE_COMPILATION ON) +endif() + diff --git a/ops/common/metax/README.md b/ops/common/metax/README.md new file mode 100644 index 0000000..19df444 --- /dev/null +++ b/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 + diff --git a/ops/common/nvidia/elementwise.cuh b/ops/common/nvidia/elementwise.cuh new file mode 100644 index 0000000..799a5d8 --- /dev/null +++ b/ops/common/nvidia/elementwise.cuh @@ -0,0 +1,45 @@ +#pragma once + +#include +#include +#include + +namespace oprt::nvidia { + +template +__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 +__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 + __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 + diff --git a/ops/common/nvidia/indexing.cuh b/ops/common/nvidia/indexing.cuh new file mode 100644 index 0000000..69c340a --- /dev/null +++ b/ops/common/nvidia/indexing.cuh @@ -0,0 +1,12 @@ +#pragma once + +#include + +namespace oprt::nvidia { + +__device__ inline int64_t contiguous_offset(int64_t linear) { + return linear; +} + +} // namespace oprt::nvidia + diff --git a/ops/common/status.cc b/ops/common/status.cc new file mode 100644 index 0000000..3c0d51d --- /dev/null +++ b/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"; + } +} + diff --git a/ops/copy/metax/README.md b/ops/copy/metax/README.md new file mode 100644 index 0000000..5fffc4b --- /dev/null +++ b/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. + diff --git a/ops/copy/metax/copy_metax.h b/ops/copy/metax/copy_metax.h new file mode 100644 index 0000000..83911f5 --- /dev/null +++ b/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. + diff --git a/ops/copy/nvidia/copy_cuda.cu b/ops/copy/nvidia/copy_cuda.cu new file mode 100644 index 0000000..4d42b5e --- /dev/null +++ b/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 + +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 +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<<>>( + static_cast(dst), static_cast(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(desc); + switch (typed->dst_view.dtype) { + case OPRT_DTYPE_F16: + return launch_copy(typed, dst, src, stream); + case OPRT_DTYPE_F32: + return launch_copy(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; +} diff --git a/ops/copy/nvidia/copy_cuda.h b/ops/copy/nvidia/copy_cuda.h new file mode 100644 index 0000000..00be5ea --- /dev/null +++ b/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 + diff --git a/ops/copy/operator.yaml b/ops/copy/operator.yaml new file mode 100644 index 0000000..4405960 --- /dev/null +++ b/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" + diff --git a/ops/copy/tilelang/copy_tl.py b/ops/copy/tilelang/copy_tl.py new file mode 100644 index 0000000..bcd9c78 --- /dev/null +++ b/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) diff --git a/ops/reduce_sum/metax/README.md b/ops/reduce_sum/metax/README.md new file mode 100644 index 0000000..557cba3 --- /dev/null +++ b/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. + diff --git a/ops/reduce_sum/metax/reduce_sum_metax.h b/ops/reduce_sum/metax/reduce_sum_metax.h new file mode 100644 index 0000000..83911f5 --- /dev/null +++ b/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. + diff --git a/ops/reduce_sum/nvidia/reduce_sum_cuda.cu b/ops/reduce_sum/nvidia/reduce_sum_cuda.cu new file mode 100644 index 0000000..f150391 --- /dev/null +++ b/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 +#include + +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(desc); + constexpr int threads = 256; + reduce_sum_rowwise_kernel<<rows, threads, threads * sizeof(float), oprt::as_cuda_stream(stream)>>>( + static_cast(out), static_cast(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; +} + diff --git a/ops/reduce_sum/nvidia/reduce_sum_cuda.h b/ops/reduce_sum/nvidia/reduce_sum_cuda.h new file mode 100644 index 0000000..303bed5 --- /dev/null +++ b/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 + diff --git a/ops/reduce_sum/operator.yaml b/ops/reduce_sum/operator.yaml new file mode 100644 index 0000000..31d2c52 --- /dev/null +++ b/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" + diff --git a/ops/reduce_sum/tilelang/reduce_sum_tl.py b/ops/reduce_sum/tilelang/reduce_sum_tl.py new file mode 100644 index 0000000..eb8e796 --- /dev/null +++ b/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) diff --git a/ops/softmax/metax/README.md b/ops/softmax/metax/README.md new file mode 100644 index 0000000..66dc7d3 --- /dev/null +++ b/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. + diff --git a/ops/softmax/metax/softmax_metax.h b/ops/softmax/metax/softmax_metax.h new file mode 100644 index 0000000..83911f5 --- /dev/null +++ b/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. + diff --git a/ops/softmax/nvidia/softmax_cuda.cu b/ops/softmax/nvidia/softmax_cuda.cu new file mode 100644 index 0000000..00479af --- /dev/null +++ b/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 +#include +#include + +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(desc); + constexpr int threads = 256; + softmax_rowwise_kernel<<rows, threads, threads * sizeof(float), oprt::as_cuda_stream(stream)>>>( + static_cast(out), static_cast(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; +} + diff --git a/ops/softmax/nvidia/softmax_cuda.h b/ops/softmax/nvidia/softmax_cuda.h new file mode 100644 index 0000000..f4bc6a4 --- /dev/null +++ b/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 + diff --git a/ops/softmax/operator.yaml b/ops/softmax/operator.yaml new file mode 100644 index 0000000..a89082a --- /dev/null +++ b/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" + diff --git a/ops/softmax/tilelang/softmax_tl.py b/ops/softmax/tilelang/softmax_tl.py new file mode 100644 index 0000000..3a347c6 --- /dev/null +++ b/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) diff --git a/ops/vector_add/metax/README.md b/ops/vector_add/metax/README.md new file mode 100644 index 0000000..a6211d3 --- /dev/null +++ b/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. + diff --git a/ops/vector_add/metax/vector_add_metax.h b/ops/vector_add/metax/vector_add_metax.h new file mode 100644 index 0000000..83911f5 --- /dev/null +++ b/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. + diff --git a/ops/vector_add/nvidia/vector_add_cuda.cu b/ops/vector_add/nvidia/vector_add_cuda.cu new file mode 100644 index 0000000..94cc52c --- /dev/null +++ b/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 + +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 +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<<>>( + static_cast(out), static_cast(a), static_cast(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(desc); + switch (typed->out_view.dtype) { + case OPRT_DTYPE_F16: + return launch_vector_add(typed, out, a, b, stream); + case OPRT_DTYPE_F32: + return launch_vector_add(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; +} diff --git a/ops/vector_add/nvidia/vector_add_cuda.h b/ops/vector_add/nvidia/vector_add_cuda.h new file mode 100644 index 0000000..0e912df --- /dev/null +++ b/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 + diff --git a/ops/vector_add/operator.yaml b/ops/vector_add/operator.yaml new file mode 100644 index 0000000..69e7ed7 --- /dev/null +++ b/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" + diff --git a/ops/vector_add/tilelang/vector_add_tl.py b/ops/vector_add/tilelang/vector_add_tl.py new file mode 100644 index 0000000..272798f --- /dev/null +++ b/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) diff --git a/python/operator_runtime/__init__.py b/python/operator_runtime/__init__.py new file mode 100644 index 0000000..ac51fa1 --- /dev/null +++ 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", +] + diff --git a/python/operator_runtime/backend.py b/python/operator_runtime/backend.py new file mode 100644 index 0000000..98bd9b2 --- /dev/null +++ b/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 + diff --git a/python/operator_runtime/benchmark.py b/python/operator_runtime/benchmark.py new file mode 100644 index 0000000..2371741 --- /dev/null +++ b/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 + diff --git a/python/operator_runtime/ctypes_bindings.py b/python/operator_runtime/ctypes_bindings.py new file mode 100644 index 0000000..1297190 --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/generated/operator_registry.py b/python/operator_runtime/generated/operator_registry.py new file mode 100644 index 0000000..37d4af1 --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/loader.py b/python/operator_runtime/loader.py new file mode 100644 index 0000000..c894794 --- /dev/null +++ b/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}") + diff --git a/python/operator_runtime/ops/copy.py b/python/operator_runtime/ops/copy.py new file mode 100644 index 0000000..bcd7977 --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/ops/reduce_sum.py b/python/operator_runtime/ops/reduce_sum.py new file mode 100644 index 0000000..fa87f7e --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/ops/softmax.py b/python/operator_runtime/ops/softmax.py new file mode 100644 index 0000000..bbe731a --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/ops/vector_add.py b/python/operator_runtime/ops/vector_add.py new file mode 100644 index 0000000..4ca4fc1 --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/perf_model.py b/python/operator_runtime/perf_model.py new file mode 100644 index 0000000..ba46396 --- /dev/null +++ b/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() + diff --git a/python/operator_runtime/prepared.py b/python/operator_runtime/prepared.py new file mode 100644 index 0000000..05e7119 --- /dev/null +++ b/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 + diff --git a/python/operator_runtime/profiler.py b/python/operator_runtime/profiler.py new file mode 100644 index 0000000..fabe7ac --- /dev/null +++ b/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 + diff --git a/python/operator_runtime/tensor_view.py b/python/operator_runtime/tensor_view.py new file mode 100644 index 0000000..d2a6eb4 --- /dev/null +++ b/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) + diff --git a/python/operator_runtime/testing.py b/python/operator_runtime/testing.py new file mode 100644 index 0000000..fac26b2 --- /dev/null +++ b/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) + diff --git a/tools/generate_operator_artifacts.py b/tools/generate_operator_artifacts.py new file mode 100644 index 0000000..d531733 --- /dev/null +++ b/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" $<$:${{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()) + diff --git a/tools/operator_schema.yaml b/tools/operator_schema.yaml new file mode 100644 index 0000000..7fbea71 --- /dev/null +++ b/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 + diff --git a/tools/validate_operator_manifest.py b/tools/validate_operator_manifest.py new file mode 100644 index 0000000..497dc0a --- /dev/null +++ b/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()) +