AddingZero
/
Intro-ops
forked from ccf-ai-infra/Intro-ops
0
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

refactor: split operator headers into public and detail layers

This commit is contained in:
yutianyu 2026-05-08 11:18:47 +08:00
parent 6e02c3eb55
commit 757dbbec74
25 changed files with 125 additions and 261 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
README.md
View File

@ -69,7 +69,7 @@ The TileLang backend requires the `tilelang` Python package. The MetaX backend i
| Training concept | Production equivalent |
| --- | --- |
| directory convention `ops/<op>/nvidia/*.cu` | build system auto-discovery / operator registry |
| C header `ops/<op>/nvidia/<op>_cuda.h` | reviewed operator API contract |
| C header `include/operator_runtime/ops/<op>.h` | reviewed operator API contract |
| descriptor lifecycle | create, workspace, execute, destroy |
| `tests/cases/<op>.py` | correctness, layout, and API contract coverage |
| `PerformanceResult` | profiler report row with latency, bytes, flops, bandwidth |

2
README.zh.md
View File

@ -68,7 +68,7 @@ TileLang 后端需要安装 `tilelang` Python 包。MetaX 后端使用独立构
| 训练概念 | 生产等价物 |
| --- | --- |
| `ops/<op>/nvidia/*.cu` 目录约定 | 构建系统自动发现 / 算子注册 |
| `ops/<op>/nvidia/<op>_cuda.h` 头文件 | 经过评审的算子 API 契约 |
| `include/operator_runtime/ops/<op>.h` 头文件 | 经过评审的算子 API 契约 |
| descriptor 生命周期 | create、workspace、execute、destroy |
| `tests/cases/<op>.py` | 正确性、布局和 API 契约覆盖 |
| `PerformanceResult` | 包含延迟、字节数、FLOPs、带宽的 profiler 报表行 |

13
include/operator_runtime/cuda_helpers.h → include/operator_runtime/detail/cuda_helpers.h
View File

@ -5,12 +5,12 @@
#ifdef __CUDACC__
#include <cuda_runtime.h>
#define OPRT_CUDA_RETURN_IF_ERROR(expr) \
do { \
cudaError_t err__ = (expr); \
if (err__ != cudaSuccess) { \
return OPRT_ERR_RUNTIME; \
} \
#define OPRT_CUDA_RETURN_IF_ERROR(expr) \
do { \
cudaError_t err__ = (expr); \
if (err__ != cudaSuccess) { \
return OPRT_ERR_RUNTIME; \
} \
} while (0)
namespace oprt {
@ -27,4 +27,3 @@ inline int blocks_for(int64_t n, int threads) {
} // namespace oprt
#endif

1
include/operator_runtime/elementwise.h → include/operator_runtime/detail/elementwise.h
View File

@ -21,4 +21,3 @@ inline bool elementwise_fast_path(const oprt_tensor_view_t &out,
} // namespace oprt
#endif

5
include/operator_runtime/operation.h → include/operator_runtime/detail/operation.h
View File

@ -1,7 +1,5 @@
#pragma once
#ifdef __cplusplus
#include <string>
namespace oprt {
@ -13,6 +11,3 @@ struct OperationSpec {
};
} // namespace oprt
#endif

1
include/operator_runtime/tensor_checks.h → include/operator_runtime/detail/tensor_checks.h
View File

@ -37,4 +37,3 @@ inline oprt_status_t check_same_shape(const oprt_tensor_view_t &a,
} // namespace oprt
#endif

9
include/operator_runtime/operator_runtime.h Normal file
View File

@ -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"

8
ops/copy/metax/copy_metax.h → include/operator_runtime/ops/copy.h
View File

@ -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

8
ops/softmax/metax/softmax_metax.h → include/operator_runtime/ops/reduce_sum.h
View File

@ -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

9
ops/softmax/nvidia/softmax_cuda.h → include/operator_runtime/ops/softmax.h
View File

@ -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

8
ops/vector_add/metax/vector_add_metax.h → include/operator_runtime/ops/vector_add.h
View File

@ -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

8
ops/CMakeLists.txt
View File

@ -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)

16
ops/copy/metax/copy_metax.maca
View File

@ -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 <cuda_fp16.h>
@ -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;

16
ops/copy/nvidia/copy_cuda.cu
View File

@ -1,9 +1,9 @@
#include "ops/copy/nvidia/copy_cuda.h"
#include "operator_runtime/ops/copy.h"
#include "operator_runtime/descriptor.h"
#include "operator_runtime/elementwise.h"
#include "operator_runtime/tensor_checks.h"
#include "operator_runtime/cuda_helpers.h"
#include "operator_runtime/detail/elementwise.h"
#include "operator_runtime/detail/tensor_checks.h"
#include "operator_runtime/detail/cuda_helpers.h"
#include "ops/copy/nvidia/kernel.cuh"
#include <cuda_fp16.h>
@ -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;

32
ops/copy/nvidia/copy_cuda.h
View File

@ -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

32
ops/reduce_sum/metax/reduce_sum_metax.h
View File

@ -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

14
ops/reduce_sum/metax/reduce_sum_metax.maca
View File

@ -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 <cuda_runtime.h>
@ -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;

14
ops/reduce_sum/nvidia/reduce_sum_cuda.cu
View File

@ -1,8 +1,8 @@
#include "ops/reduce_sum/nvidia/reduce_sum_cuda.h"
#include "operator_runtime/ops/reduce_sum.h"
#include "operator_runtime/descriptor.h"
#include "operator_runtime/tensor_checks.h"
#include "operator_runtime/cuda_helpers.h"
#include "operator_runtime/detail/tensor_checks.h"
#include "operator_runtime/detail/cuda_helpers.h"
#include "ops/reduce_sum/nvidia/kernel.cuh"
#include <cuda_runtime.h>
@ -34,7 +34,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in
} // namespace
extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia(
extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor(
oprt_operator_descriptor_t *desc,
const oprt_tensor_view_t *out,
const oprt_tensor_view_t *in,
@ -66,7 +66,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_reduce_sum_descriptor_nvidia(
return OPRT_SUCCESS;
}
extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia(
extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size(
oprt_operator_descriptor_t desc,
size_t *size) {
if (desc == nullptr || size == nullptr) {
@ -76,7 +76,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_reduce_sum_workspace_size_nvidia(
return OPRT_SUCCESS;
}
extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum_nvidia(
extern "C" OPRT_EXPORT oprt_status_t oprt_execute_reduce_sum(
oprt_operator_descriptor_t desc,
void *,
size_t workspace_size,
@ -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;

33
ops/reduce_sum/nvidia/reduce_sum_cuda.h
View File

@ -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

14
ops/softmax/metax/softmax_metax.maca
View File

@ -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 <cuda_runtime.h>
#include <float.h>
@ -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;

14
ops/softmax/nvidia/softmax_cuda.cu
View File

@ -1,8 +1,8 @@
#include "ops/softmax/nvidia/softmax_cuda.h"
#include "operator_runtime/ops/softmax.h"
#include "operator_runtime/descriptor.h"
#include "operator_runtime/tensor_checks.h"
#include "operator_runtime/cuda_helpers.h"
#include "operator_runtime/detail/tensor_checks.h"
#include "operator_runtime/detail/cuda_helpers.h"
#include "ops/softmax/nvidia/kernel.cuh"
#include <cuda_runtime.h>
@ -34,7 +34,7 @@ bool is_rowwise_case(const oprt_tensor_view_t &out, const oprt_tensor_view_t &in
} // namespace
extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia(
extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor(
oprt_operator_descriptor_t *desc,
const oprt_tensor_view_t *out,
const oprt_tensor_view_t *in,
@ -66,7 +66,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_create_softmax_descriptor_nvidia(
return OPRT_SUCCESS;
}
extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia(
extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size(
oprt_operator_descriptor_t desc,
size_t *size) {
if (desc == nullptr || size == nullptr) {
@ -76,7 +76,7 @@ extern "C" OPRT_EXPORT oprt_status_t oprt_get_softmax_workspace_size_nvidia(
return OPRT_SUCCESS;
}
extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax_nvidia(
extern "C" OPRT_EXPORT oprt_status_t oprt_execute_softmax(
oprt_operator_descriptor_t desc,
void *,
size_t workspace_size,
@ -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;

16
ops/vector_add/metax/vector_add_metax.maca
View File

@ -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 <cuda_fp16.h>
@ -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;

16
ops/vector_add/nvidia/vector_add_cuda.cu
View File

@ -1,9 +1,9 @@
#include "ops/vector_add/nvidia/vector_add_cuda.h"
#include "operator_runtime/ops/vector_add.h"
#include "operator_runtime/descriptor.h"
#include "operator_runtime/elementwise.h"
#include "operator_runtime/tensor_checks.h"
#include "operator_runtime/cuda_helpers.h"
#include "operator_runtime/detail/elementwise.h"
#include "operator_runtime/detail/tensor_checks.h"
#include "operator_runtime/detail/cuda_helpers.h"
#include "ops/vector_add/nvidia/kernel.cuh"
#include <cuda_fp16.h>
@ -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;

34
ops/vector_add/nvidia/vector_add_cuda.h
View File

@ -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

61
python/operator_runtime/_internal/bindings.py
View File

@ -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)