From 757dbbec74e247a954d6937190a0322aa7e6e676 Mon Sep 17 00:00:00 2001 From: yutianyu Date: Fri, 8 May 2026 11:18:47 +0800 Subject: [PATCH] refactor: split operator headers into public and detail layers --- README.md | 2 +- README.zh.md | 2 +- .../{ => detail}/cuda_helpers.h | 13 ++-- .../{ => detail}/elementwise.h | 1 - .../operator_runtime/{ => detail}/operation.h | 5 -- .../{ => detail}/tensor_checks.h | 1 - include/operator_runtime/operator_runtime.h | 9 +++ .../operator_runtime/ops/copy.h | 8 +-- .../operator_runtime/ops/reduce_sum.h | 8 +-- .../operator_runtime/ops/softmax.h | 9 ++- .../operator_runtime/ops/vector_add.h | 8 +-- ops/CMakeLists.txt | 8 ++- ops/copy/metax/copy_metax.maca | 16 ++--- ops/copy/nvidia/copy_cuda.cu | 16 ++--- ops/copy/nvidia/copy_cuda.h | 32 ---------- ops/reduce_sum/metax/reduce_sum_metax.h | 32 ---------- ops/reduce_sum/metax/reduce_sum_metax.maca | 14 ++--- ops/reduce_sum/nvidia/reduce_sum_cuda.cu | 14 ++--- ops/reduce_sum/nvidia/reduce_sum_cuda.h | 33 ---------- ops/softmax/metax/softmax_metax.maca | 14 ++--- ops/softmax/nvidia/softmax_cuda.cu | 14 ++--- ops/vector_add/metax/vector_add_metax.maca | 16 ++--- ops/vector_add/nvidia/vector_add_cuda.cu | 16 ++--- ops/vector_add/nvidia/vector_add_cuda.h | 34 ----------- python/operator_runtime/_internal/bindings.py | 61 ++++++++----------- 25 files changed, 125 insertions(+), 261 deletions(-) rename include/operator_runtime/{ => detail}/cuda_helpers.h (59%) rename include/operator_runtime/{ => detail}/elementwise.h (99%) rename include/operator_runtime/{ => detail}/operation.h (85%) rename include/operator_runtime/{ => detail}/tensor_checks.h (99%) create mode 100644 include/operator_runtime/operator_runtime.h rename ops/copy/metax/copy_metax.h => include/operator_runtime/ops/copy.h (66%) rename ops/softmax/metax/softmax_metax.h => include/operator_runtime/ops/reduce_sum.h (66%) rename ops/softmax/nvidia/softmax_cuda.h => include/operator_runtime/ops/softmax.h (65%) rename ops/vector_add/metax/vector_add_metax.h => include/operator_runtime/ops/vector_add.h (66%) delete mode 100644 ops/copy/nvidia/copy_cuda.h delete mode 100644 ops/reduce_sum/metax/reduce_sum_metax.h delete mode 100644 ops/reduce_sum/nvidia/reduce_sum_cuda.h delete mode 100644 ops/vector_add/nvidia/vector_add_cuda.h diff --git a/README.md b/README.md index 02a5a9a..e800c52 100644 --- a/README.md +++ b/README.md @@ -69,7 +69,7 @@ The TileLang backend requires the `tilelang` Python package. The MetaX backend i | Training concept | Production equivalent | | --- | --- | | directory convention `ops//nvidia/*.cu` | build system auto-discovery / operator registry | -| C header `ops//nvidia/_cuda.h` | reviewed operator API contract | +| C header `include/operator_runtime/ops/.h` | reviewed operator API contract | | descriptor lifecycle | create, workspace, execute, destroy | | `tests/cases/.py` | correctness, layout, and API contract coverage | | `PerformanceResult` | profiler report row with latency, bytes, flops, bandwidth | diff --git a/README.zh.md b/README.zh.md index d168e29..3a86606 100644 --- a/README.zh.md +++ b/README.zh.md @@ -68,7 +68,7 @@ TileLang 后端需要安装 `tilelang` Python 包。MetaX 后端使用独立构 | 训练概念 | 生产等价物 | | --- | --- | | `ops//nvidia/*.cu` 目录约定 | 构建系统自动发现 / 算子注册 | -| `ops//nvidia/_cuda.h` 头文件 | 经过评审的算子 API 契约 | +| `include/operator_runtime/ops/.h` 头文件 | 经过评审的算子 API 契约 | | descriptor 生命周期 | create、workspace、execute、destroy | | `tests/cases/.py` | 正确性、布局和 API 契约覆盖 | | `PerformanceResult` | 包含延迟、字节数、FLOPs、带宽的 profiler 报表行 | diff --git a/include/operator_runtime/cuda_helpers.h b/include/operator_runtime/detail/cuda_helpers.h similarity index 59% rename from include/operator_runtime/cuda_helpers.h rename to include/operator_runtime/detail/cuda_helpers.h index 1fbac5d..d155529 100644 --- a/include/operator_runtime/cuda_helpers.h +++ b/include/operator_runtime/detail/cuda_helpers.h @@ -5,12 +5,12 @@ #ifdef __CUDACC__ #include -#define OPRT_CUDA_RETURN_IF_ERROR(expr) \ - do { \ - cudaError_t err__ = (expr); \ - if (err__ != cudaSuccess) { \ - return OPRT_ERR_RUNTIME; \ - } \ +#define OPRT_CUDA_RETURN_IF_ERROR(expr) \ + do { \ + cudaError_t err__ = (expr); \ + if (err__ != cudaSuccess) { \ + return OPRT_ERR_RUNTIME; \ + } \ } while (0) namespace oprt { @@ -27,4 +27,3 @@ inline int blocks_for(int64_t n, int threads) { } // namespace oprt #endif - diff --git a/include/operator_runtime/elementwise.h b/include/operator_runtime/detail/elementwise.h similarity index 99% rename from include/operator_runtime/elementwise.h rename to include/operator_runtime/detail/elementwise.h index f7d42ed..653787f 100644 --- a/include/operator_runtime/elementwise.h +++ b/include/operator_runtime/detail/elementwise.h @@ -21,4 +21,3 @@ inline bool elementwise_fast_path(const oprt_tensor_view_t &out, } // namespace oprt #endif - diff --git a/include/operator_runtime/operation.h b/include/operator_runtime/detail/operation.h similarity index 85% rename from include/operator_runtime/operation.h rename to include/operator_runtime/detail/operation.h index 88a61e2..2446529 100644 --- a/include/operator_runtime/operation.h +++ b/include/operator_runtime/detail/operation.h @@ -1,7 +1,5 @@ #pragma once -#ifdef __cplusplus - #include namespace oprt { @@ -13,6 +11,3 @@ struct OperationSpec { }; } // namespace oprt - -#endif - diff --git a/include/operator_runtime/tensor_checks.h b/include/operator_runtime/detail/tensor_checks.h similarity index 99% rename from include/operator_runtime/tensor_checks.h rename to include/operator_runtime/detail/tensor_checks.h index 47063df..0f794b0 100644 --- a/include/operator_runtime/tensor_checks.h +++ b/include/operator_runtime/detail/tensor_checks.h @@ -37,4 +37,3 @@ inline oprt_status_t check_same_shape(const oprt_tensor_view_t &a, } // namespace oprt #endif - diff --git a/include/operator_runtime/operator_runtime.h b/include/operator_runtime/operator_runtime.h new file mode 100644 index 0000000..06b06e7 --- /dev/null +++ b/include/operator_runtime/operator_runtime.h @@ -0,0 +1,9 @@ +#pragma once + +#include "operator_runtime/api.h" +#include "operator_runtime/descriptor.h" +#include "operator_runtime/tensor_view.h" +#include "operator_runtime/ops/copy.h" +#include "operator_runtime/ops/reduce_sum.h" +#include "operator_runtime/ops/softmax.h" +#include "operator_runtime/ops/vector_add.h" diff --git a/ops/copy/metax/copy_metax.h b/include/operator_runtime/ops/copy.h similarity index 66% rename from ops/copy/metax/copy_metax.h rename to include/operator_runtime/ops/copy.h index 092ed3f..73e8138 100644 --- a/ops/copy/metax/copy_metax.h +++ b/include/operator_runtime/ops/copy.h @@ -6,16 +6,16 @@ extern "C" { #endif -OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_metax( +OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *dst, const oprt_tensor_view_t *src); -OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_metax( +OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size( oprt_operator_descriptor_t desc, size_t *size); -OPRT_EXPORT oprt_status_t oprt_execute_copy_metax( +OPRT_EXPORT oprt_status_t oprt_execute_copy( oprt_operator_descriptor_t desc, void *workspace, size_t workspace_size, @@ -23,7 +23,7 @@ OPRT_EXPORT oprt_status_t oprt_execute_copy_metax( const void *src, oprt_stream_t stream); -OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor_metax( +OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor( oprt_operator_descriptor_t desc); #ifdef __cplusplus diff --git a/ops/softmax/metax/softmax_metax.h b/include/operator_runtime/ops/reduce_sum.h similarity index 66% rename from ops/softmax/metax/softmax_metax.h rename to include/operator_runtime/ops/reduce_sum.h index 103c8af..eca223d 100644 --- a/ops/softmax/metax/softmax_metax.h +++ b/include/operator_runtime/ops/reduce_sum.h @@ -6,17 +6,17 @@ extern "C" { #endif -OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_metax( +OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *in, int64_t axis); -OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_metax( +OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size( oprt_operator_descriptor_t desc, size_t *size); -OPRT_EXPORT oprt_status_t oprt_execute_softmax_metax( +OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum( oprt_operator_descriptor_t desc, void *workspace, size_t workspace_size, @@ -24,7 +24,7 @@ OPRT_EXPORT oprt_status_t oprt_execute_softmax_metax( const void *in, oprt_stream_t stream); -OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_metax( +OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor( oprt_operator_descriptor_t desc); #ifdef __cplusplus diff --git a/ops/softmax/nvidia/softmax_cuda.h b/include/operator_runtime/ops/softmax.h similarity index 65% rename from ops/softmax/nvidia/softmax_cuda.h rename to include/operator_runtime/ops/softmax.h index f4bc6a4..2b929ff 100644 --- a/ops/softmax/nvidia/softmax_cuda.h +++ b/include/operator_runtime/ops/softmax.h @@ -6,17 +6,17 @@ extern "C" { #endif -OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia( +OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *in, int64_t axis); -OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia( +OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size( oprt_operator_descriptor_t desc, size_t *size); -OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia( +OPRT_EXPORT oprt_status_t oprt_execute_softmax( oprt_operator_descriptor_t desc, void *workspace, size_t workspace_size, @@ -24,10 +24,9 @@ OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia( const void *in, oprt_stream_t stream); -OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_nvidia( +OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor( oprt_operator_descriptor_t desc); #ifdef __cplusplus } #endif - diff --git a/ops/vector_add/metax/vector_add_metax.h b/include/operator_runtime/ops/vector_add.h similarity index 66% rename from ops/vector_add/metax/vector_add_metax.h rename to include/operator_runtime/ops/vector_add.h index 75a035d..85af267 100644 --- a/ops/vector_add/metax/vector_add_metax.h +++ b/include/operator_runtime/ops/vector_add.h @@ -6,17 +6,17 @@ extern "C" { #endif -OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_metax( +OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *a, const oprt_tensor_view_t *b); -OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_metax( +OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size( oprt_operator_descriptor_t desc, size_t *size); -OPRT_EXPORT oprt_status_t oprt_execute_vector_add_metax( +OPRT_EXPORT oprt_status_t oprt_execute_vector_add( oprt_operator_descriptor_t desc, void *workspace, size_t workspace_size, @@ -25,7 +25,7 @@ OPRT_EXPORT oprt_status_t oprt_execute_vector_add_metax( const void *b, oprt_stream_t stream); -OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor_metax( +OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor( oprt_operator_descriptor_t desc); #ifdef __cplusplus diff --git a/ops/CMakeLists.txt b/ops/CMakeLists.txt index 25b8af1..80ca0eb 100644 --- a/ops/CMakeLists.txt +++ b/ops/CMakeLists.txt @@ -25,9 +25,11 @@ endif() set(CAMP_OPERATOR_SOURCES ${CAMP_COMMON_SOURCES} ${CAMP_NVIDIA_SOURCES} ${CAMP_METAX_SOURCES}) add_library(camp_ops SHARED ${CAMP_OPERATOR_SOURCES}) -target_include_directories(camp_ops PUBLIC - "${CMAKE_CURRENT_SOURCE_DIR}/.." - "${CMAKE_CURRENT_SOURCE_DIR}/../include" +target_include_directories(camp_ops + PUBLIC + "${CMAKE_CURRENT_SOURCE_DIR}/../include" + PRIVATE + "${CMAKE_CURRENT_SOURCE_DIR}/.." ) if(CAMP_ENABLE_METAX) diff --git a/ops/copy/metax/copy_metax.maca b/ops/copy/metax/copy_metax.maca index 367cde8..832a445 100644 --- a/ops/copy/metax/copy_metax.maca +++ b/ops/copy/metax/copy_metax.maca @@ -1,9 +1,9 @@ -#include "ops/copy/metax/copy_metax.h" +#include "operator_runtime/ops/copy.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/elementwise.h" -#include "operator_runtime/tensor_checks.h" +#include "operator_runtime/detail/elementwise.h" +#include "operator_runtime/detail/tensor_checks.h" #include @@ -41,7 +41,7 @@ oprt_status_t launch_copy(const CopyDescriptor *desc, void *dst, const void *src } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *dst, const oprt_tensor_view_t *src) { @@ -78,7 +78,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -88,7 +88,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -112,7 +112,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy_metax( } } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/copy/nvidia/copy_cuda.cu b/ops/copy/nvidia/copy_cuda.cu index d5d8f0f..c054ea2 100644 --- a/ops/copy/nvidia/copy_cuda.cu +++ b/ops/copy/nvidia/copy_cuda.cu @@ -1,9 +1,9 @@ -#include "ops/copy/nvidia/copy_cuda.h" +#include "operator_runtime/ops/copy.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/elementwise.h" -#include "operator_runtime/tensor_checks.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/elementwise.h" +#include "operator_runtime/detail/tensor_checks.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "ops/copy/nvidia/kernel.cuh" #include @@ -33,7 +33,7 @@ oprt_status_t launch_copy(const CopyDescriptor *desc, void *dst, const void *src } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *dst, const oprt_tensor_view_t *src) { @@ -70,7 +70,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_copy_descriptor_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -80,7 +80,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_copy_workspace_size_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -104,7 +104,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_copy_nvidia( } } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_copy_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/copy/nvidia/copy_cuda.h b/ops/copy/nvidia/copy_cuda.h deleted file mode 100644 index 00be5ea..0000000 --- a/ops/copy/nvidia/copy_cuda.h +++ /dev/null @@ -1,32 +0,0 @@ -#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/reduce_sum/metax/reduce_sum_metax.h b/ops/reduce_sum/metax/reduce_sum_metax.h deleted file mode 100644 index fc7288a..0000000 --- a/ops/reduce_sum/metax/reduce_sum_metax.h +++ /dev/null @@ -1,32 +0,0 @@ -#pragma once - -#include "operator_runtime/api.h" - -#ifdef __cplusplus -extern "C" { -#endif - -OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_metax( - 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_metax( - oprt_operator_descriptor_t desc, - size_t *size); - -OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_metax( - 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_metax( - oprt_operator_descriptor_t desc); - -#ifdef __cplusplus -} -#endif diff --git a/ops/reduce_sum/metax/reduce_sum_metax.maca b/ops/reduce_sum/metax/reduce_sum_metax.maca index e7b88ab..5ea6c72 100644 --- a/ops/reduce_sum/metax/reduce_sum_metax.maca +++ b/ops/reduce_sum/metax/reduce_sum_metax.maca @@ -1,8 +1,8 @@ -#include "ops/reduce_sum/metax/reduce_sum_metax.h" +#include "operator_runtime/ops/reduce_sum.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/tensor_checks.h" +#include "operator_runtime/detail/tensor_checks.h" #include @@ -61,7 +61,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *in, @@ -93,7 +93,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -103,7 +103,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -125,7 +125,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/reduce_sum/nvidia/reduce_sum_cuda.cu b/ops/reduce_sum/nvidia/reduce_sum_cuda.cu index c20418e..a383af1 100644 --- a/ops/reduce_sum/nvidia/reduce_sum_cuda.cu +++ b/ops/reduce_sum/nvidia/reduce_sum_cuda.cu @@ -1,8 +1,8 @@ -#include "ops/reduce_sum/nvidia/reduce_sum_cuda.h" +#include "operator_runtime/ops/reduce_sum.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/tensor_checks.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/tensor_checks.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "ops/reduce_sum/nvidia/kernel.cuh" #include @@ -34,7 +34,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *in, @@ -66,7 +66,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -76,7 +76,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -98,7 +98,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_reduce_sum_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/reduce_sum/nvidia/reduce_sum_cuda.h b/ops/reduce_sum/nvidia/reduce_sum_cuda.h deleted file mode 100644 index 303bed5..0000000 --- a/ops/reduce_sum/nvidia/reduce_sum_cuda.h +++ /dev/null @@ -1,33 +0,0 @@ -#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/softmax/metax/softmax_metax.maca b/ops/softmax/metax/softmax_metax.maca index 6c3c692..d4167b5 100644 --- a/ops/softmax/metax/softmax_metax.maca +++ b/ops/softmax/metax/softmax_metax.maca @@ -1,8 +1,8 @@ -#include "ops/softmax/metax/softmax_metax.h" +#include "operator_runtime/ops/softmax.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/tensor_checks.h" +#include "operator_runtime/detail/tensor_checks.h" #include #include @@ -84,7 +84,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *in, @@ -116,7 +116,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -126,7 +126,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -148,7 +148,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/softmax/nvidia/softmax_cuda.cu b/ops/softmax/nvidia/softmax_cuda.cu index 1c561b0..5afeb4f 100644 --- a/ops/softmax/nvidia/softmax_cuda.cu +++ b/ops/softmax/nvidia/softmax_cuda.cu @@ -1,8 +1,8 @@ -#include "ops/softmax/nvidia/softmax_cuda.h" +#include "operator_runtime/ops/softmax.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/tensor_checks.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/tensor_checks.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "ops/softmax/nvidia/kernel.cuh" #include @@ -34,7 +34,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *in, @@ -66,7 +66,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -76,7 +76,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -98,7 +98,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_softmax_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/vector_add/metax/vector_add_metax.maca b/ops/vector_add/metax/vector_add_metax.maca index 5279742..1122f10 100644 --- a/ops/vector_add/metax/vector_add_metax.maca +++ b/ops/vector_add/metax/vector_add_metax.maca @@ -1,9 +1,9 @@ -#include "ops/vector_add/metax/vector_add_metax.h" +#include "operator_runtime/ops/vector_add.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/elementwise.h" -#include "operator_runtime/tensor_checks.h" +#include "operator_runtime/detail/elementwise.h" +#include "operator_runtime/detail/tensor_checks.h" #include @@ -56,7 +56,7 @@ oprt_status_t launch_vector_add(const VectorAddDescriptor *desc, void *out, cons } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *a, @@ -99,7 +99,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -109,7 +109,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_metax( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -134,7 +134,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add_metax( } } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor_metax( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/vector_add/nvidia/vector_add_cuda.cu b/ops/vector_add/nvidia/vector_add_cuda.cu index b88bc4a..1b5ace2 100644 --- a/ops/vector_add/nvidia/vector_add_cuda.cu +++ b/ops/vector_add/nvidia/vector_add_cuda.cu @@ -1,9 +1,9 @@ -#include "ops/vector_add/nvidia/vector_add_cuda.h" +#include "operator_runtime/ops/vector_add.h" #include "operator_runtime/descriptor.h" -#include "operator_runtime/elementwise.h" -#include "operator_runtime/tensor_checks.h" -#include "operator_runtime/cuda_helpers.h" +#include "operator_runtime/detail/elementwise.h" +#include "operator_runtime/detail/tensor_checks.h" +#include "operator_runtime/detail/cuda_helpers.h" #include "ops/vector_add/nvidia/kernel.cuh" #include @@ -34,7 +34,7 @@ oprt_status_t launch_vector_add(const VectorAddDescriptor *desc, void *out, cons } // namespace -extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor( oprt_operator_descriptor_t *desc, const oprt_tensor_view_t *out, const oprt_tensor_view_t *a, @@ -77,7 +77,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_vector_add_descriptor_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size( oprt_operator_descriptor_t desc, size_t *size) { if (desc == nullptr || size == nullptr) { @@ -87,7 +87,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_vector_add_workspace_size_nvidia( return OPRT_SUCCESS; } -extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add( oprt_operator_descriptor_t desc, void *, size_t workspace_size, @@ -112,7 +112,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_execute_vector_add_nvidia( } } -extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor_nvidia( +extern "C" OPRT_EXPORT oprt_status_t oprt_destroy_vector_add_descriptor( oprt_operator_descriptor_t desc) { delete desc; return OPRT_SUCCESS; diff --git a/ops/vector_add/nvidia/vector_add_cuda.h b/ops/vector_add/nvidia/vector_add_cuda.h deleted file mode 100644 index 0e912df..0000000 --- a/ops/vector_add/nvidia/vector_add_cuda.h +++ /dev/null @@ -1,34 +0,0 @@ -#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/python/operator_runtime/_internal/bindings.py b/python/operator_runtime/_internal/bindings.py index 906e2ac..23b6361 100644 --- a/python/operator_runtime/_internal/bindings.py +++ b/python/operator_runtime/_internal/bindings.py @@ -4,7 +4,7 @@ import ctypes from dataclasses import dataclass from typing import Callable -from operator_runtime.backend import Backend, normalize_backend +from operator_runtime.backend import Backend from .loader import load_library from .tensor_view import TensorView @@ -35,40 +35,35 @@ class CFunctions: destroy: Callable -def _backend_suffix(backend: str | Backend) -> str: - return normalize_backend(backend).value - - -def _missing_symbol_error(name: str, backend: str, symbol: str, exc: AttributeError) -> OperatorRuntimeError: +def _missing_symbol_error(name: str, symbol: str, exc: AttributeError) -> OperatorRuntimeError: raise OperatorRuntimeError( - f"backend {backend} is unavailable in the loaded libcamp_ops.so; missing symbol {symbol}" + f"missing symbol {symbol} for operator {name}" ) from exc def bind_unary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions: lib = load_library() - suffix = _backend_suffix(backend) - create_symbol = f"oprt_create_{name}_descriptor_{suffix}" + create_symbol = f"oprt_create_{name}_descriptor" try: create = getattr(lib, create_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, create_symbol, exc) + _missing_symbol_error(name, create_symbol, exc) create.argtypes = [ctypes.POINTER(Descriptor), ctypes.POINTER(TensorView), ctypes.POINTER(TensorView)] create.restype = Status - workspace_symbol = f"oprt_get_{name}_workspace_size_{suffix}" + workspace_symbol = f"oprt_get_{name}_workspace_size" try: workspace = getattr(lib, workspace_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, workspace_symbol, exc) + _missing_symbol_error(name, workspace_symbol, exc) workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)] workspace.restype = Status - execute_symbol = f"oprt_execute_{name}_{suffix}" + execute_symbol = f"oprt_execute_{name}" try: execute = getattr(lib, execute_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, execute_symbol, exc) + _missing_symbol_error(name, execute_symbol, exc) execute.argtypes = [ Descriptor, ctypes.c_void_p, @@ -79,11 +74,11 @@ def bind_unary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions ] execute.restype = Status - destroy_symbol = f"oprt_destroy_{name}_descriptor_{suffix}" + destroy_symbol = f"oprt_destroy_{name}_descriptor" try: destroy = getattr(lib, destroy_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, destroy_symbol, exc) + _missing_symbol_error(name, destroy_symbol, exc) destroy.argtypes = [Descriptor] destroy.restype = Status return CFunctions(create, workspace, execute, destroy) @@ -91,12 +86,11 @@ def bind_unary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions def bind_binary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions: lib = load_library() - suffix = _backend_suffix(backend) - create_symbol = f"oprt_create_{name}_descriptor_{suffix}" + create_symbol = f"oprt_create_{name}_descriptor" try: create = getattr(lib, create_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, create_symbol, exc) + _missing_symbol_error(name, create_symbol, exc) create.argtypes = [ ctypes.POINTER(Descriptor), ctypes.POINTER(TensorView), @@ -105,19 +99,19 @@ def bind_binary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunction ] create.restype = Status - workspace_symbol = f"oprt_get_{name}_workspace_size_{suffix}" + workspace_symbol = f"oprt_get_{name}_workspace_size" try: workspace = getattr(lib, workspace_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, workspace_symbol, exc) + _missing_symbol_error(name, workspace_symbol, exc) workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)] workspace.restype = Status - execute_symbol = f"oprt_execute_{name}_{suffix}" + execute_symbol = f"oprt_execute_{name}" try: execute = getattr(lib, execute_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, execute_symbol, exc) + _missing_symbol_error(name, execute_symbol, exc) execute.argtypes = [ Descriptor, ctypes.c_void_p, @@ -129,11 +123,11 @@ def bind_binary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunction ] execute.restype = Status - destroy_symbol = f"oprt_destroy_{name}_descriptor_{suffix}" + destroy_symbol = f"oprt_destroy_{name}_descriptor" try: destroy = getattr(lib, destroy_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, destroy_symbol, exc) + _missing_symbol_error(name, destroy_symbol, exc) destroy.argtypes = [Descriptor] destroy.restype = Status return CFunctions(create, workspace, execute, destroy) @@ -141,12 +135,11 @@ def bind_binary(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunction def bind_reduce_like(name: str, backend: str | Backend = Backend.NVIDIA) -> CFunctions: lib = load_library() - suffix = _backend_suffix(backend) - create_symbol = f"oprt_create_{name}_descriptor_{suffix}" + create_symbol = f"oprt_create_{name}_descriptor" try: create = getattr(lib, create_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, create_symbol, exc) + _missing_symbol_error(name, create_symbol, exc) create.argtypes = [ ctypes.POINTER(Descriptor), ctypes.POINTER(TensorView), @@ -155,19 +148,19 @@ def bind_reduce_like(name: str, backend: str | Backend = Backend.NVIDIA) -> CFun ] create.restype = Status - workspace_symbol = f"oprt_get_{name}_workspace_size_{suffix}" + workspace_symbol = f"oprt_get_{name}_workspace_size" try: workspace = getattr(lib, workspace_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, workspace_symbol, exc) + _missing_symbol_error(name, workspace_symbol, exc) workspace.argtypes = [Descriptor, ctypes.POINTER(ctypes.c_size_t)] workspace.restype = Status - execute_symbol = f"oprt_execute_{name}_{suffix}" + execute_symbol = f"oprt_execute_{name}" try: execute = getattr(lib, execute_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, execute_symbol, exc) + _missing_symbol_error(name, execute_symbol, exc) execute.argtypes = [ Descriptor, ctypes.c_void_p, @@ -178,11 +171,11 @@ def bind_reduce_like(name: str, backend: str | Backend = Backend.NVIDIA) -> CFun ] execute.restype = Status - destroy_symbol = f"oprt_destroy_{name}_descriptor_{suffix}" + destroy_symbol = f"oprt_destroy_{name}_descriptor" try: destroy = getattr(lib, destroy_symbol) except AttributeError as exc: - _missing_symbol_error(name, suffix, destroy_symbol, exc) + _missing_symbol_error(name, destroy_symbol, exc) destroy.argtypes = [Descriptor] destroy.restype = Status return CFunctions(create, workspace, execute, destroy)