When PyTorch is used for model training and inference, many operators are called for computation. To integrate and deploy your custom operators in PyTorch, the following methods are available:

This section describes how to register a custom operator through torch.library and Pybind and enable PyTorch to call the operator kernel program.
Pybind is used to quickly expose C++ functions to Python for efficient API binding. However, the generated operators cannot be identified by the PyTorch operator system and do not support schema definition or graph tracing capability. Therefore, torch.compile is not supported. In contrast, torch.library provides a mechanism for in-depth integration with PyTorch's core operator system and supports operator registration, schema definition, and graph tracing, which are prerequisites for supporting torch.compile. You can choose the method that best fits your needs.
The following code uses the add_custom operator as an example to describe how to call the operator kernel program through torch.library. Only the core steps are described here. For details about the complete sample, see torch.library sample.
Implement the operator on the kernel and use <<<>>> to call the operator's kernel function to complete the specified computation. In the sample, the c10_npu::getCurrentNPUStream API is called to obtain the current NPU stream. The return type is NPUStream. For details about how to use this API, see "(beta) c10_npu::getCurrentNPUStream" in Ascend Extension for PyTorch Custom API Reference.
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 | namespace ascendc_ops { at::Tensor ascendc_add(const at::Tensor& x, const at::Tensor& y) { // Allocate runtime resources and obtain the current NPU stream by calling c10_npu::getCurrentNPUStream(). auto aclStream = c10_npu::getCurrentNPUStream().stream(false); // Allocate the output buffer on the device. at::Tensor z = at::empty_like(x); uint32_t numBlocks = 8; uint32_t totalLength = 1; for (uint32_t size : x.sizes()) { totalLength *= size; } // Call the kernel function using <<<>>> to complete the specified computation. add_custom<<<numBlocks, nullptr, aclStream>>>((uint8_t*)(x.mutable_data_ptr()), (uint8_t*)(y.mutable_data_ptr()), (uint8_t*)(z.mutable_data_ptr()), totalLength); // Copy the computation result from the device to the host and free the allocated resources. return z; } } // namespace ascendc_ops |
1 2 3 4 5 6 7 8 9 10 11 | // Register the operator with torch.library. TORCH_LIBRARY(ascendc_ops, m) { m.def("ascendc_add(Tensor x, Tensor y) -> Tensor"); } // Register the implementation with PrivateUse1 on the NPU device. TORCH_LIBRARY_IMPL(ascendc_ops, PrivateUse1, m) { m.impl("ascendc_add", TORCH_FN(ascendc_ops::ascendc_add)); } |
In add_custom_test.py, the generated custom operator library is loaded using torch.ops.load_library, the registered ascendc_add function is called, and the output on the NPU is compared with the CPU's standard addition result to verify the numerical correctness of the custom operator.
The following code uses the add_custom operator as an example to describe how to call a custom operator in a PyTorch script using Pybind. Only the core steps are described here. For details about the complete sample, see Pybind sample.
pip3 install pybind11 expecttest
Implement the operator on the kernel and use <<<>>> to call the operator's kernel function to complete the specified computation. In the sample, the c10_npu::getCurrentNPUStream API is called to obtain the current NPU stream. The return type is NPUStream. For details about how to use this API, see "(beta) c10_npu::getCurrentNPUStream" in Ascend Extension for PyTorch Custom API Reference.
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 | // Header files required for calling Pybind and PyTorch #include <pybind11/pybind11.h> #include <torch/extension.h> #include "torch_npu/csrc/core/npu/NPUStream.h" // Header file required for implementation on the kernel #include "kernel_operator.h" ... namespace ascendc_ops { at::Tensor ascendc_add(const at::Tensor& x, const at::Tensor& y) { // Allocate runtime resources and obtain the current NPU stream by calling c10_npu::getCurrentNPUStream(). auto aclStream = c10_npu::getCurrentNPUStream().stream(false); // Allocate the output buffer on the device. at::Tensor z = at::empty_like(x); uint32_t numBlocks = 8; uint32_t totalLength = 1; for (uint32_t size : x.sizes()) { totalLength *= size; } // Call the kernel function using <<<>>> to complete the specified computation. add_custom<<<numBlocks, nullptr, aclStream>>>((uint8_t*)(x.mutable_data_ptr()), (uint8_t*)(y.mutable_data_ptr()), (uint8_t*)(z.mutable_data_ptr()), totalLength); // Copy the computation result from the device to the host and free the allocated resources. return z; } } // namespace ascendc_ops |
1 2 3 4 5 | PYBIND11_MODULE(ascendc_ops, m)// ascendc_ops: module name; m: module object. { m.doc() = "add_custom pybind11 interfaces";// optional module docstring m.def("ascendc_add", &ascendc_ops::ascendc_add, "");// Bind the ascendc_add function to the Pybind module. } |