OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用 - 通过aclnn调用的方式调用AddCustom算子

OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用 通过aclnn调用的方式调用 - AddCustom算子 - 单算子API执行(aclnn)

多种算子调用方式

*开发时间使用场景调用方式运行硬件
基于Kernel直调工程(快速)单算子调用,快速验证算法逻辑ICPU_RUN_KF
内核调用符号
ACLRT_LAUNCH_KERNEL
CPU
NPU
NPU
基于自定义算子工程(标准)单算子网络/整网部署调用aclInnxxx
aclopExecuteV2
pytorch adapter
NPU
NPU
NPU

AscendCL单算子调用概述

完成自定义算子的开发部署(标准流程)后,可以通过单算子调用的方式来验证单算子的功能。单算子调用有API执行和模型执行两种方式:

单算子API执行(aclnn) (本文描述的):

基于C语言的API执行算子,直接调用单算子API接口。多用于大模型训练算子,既整网算子模型较为固定的场景。

单算子模型执行:

基于图模式执行算子,先编译算子(例如,使用ATC工具将AscendIR定义的单算子描述文件编译成算子om模型文件),再调用AscendCL接口加载算子模型,最后调用AscendCL接口执行算子。多用于搜广推,整网模型变化较大的场景。

前面的内容是实现AddCustom算子,这里是通过aclnn调用的方式调用AddCustom算子,也是基于C语言的API执行算子,直接调用单算子API接口。
所有有两部分

一部分是 之前的实现

/home/HwHiAiUser/samples/operator/AddCustomSample/FrameworkLaunch/AddCustom

在这里插入图片描述可以从json串开始生成代码骨架,然后开始填写内容。具体请看这里
前置知识 1
前置知识 2
前置知识 3

编译
在这里插入图片描述

因为单算子API执行方式,会自动在编译工程的build_out/autogen目录下生成.cpp和.h,编写单算子的调用代码时,要包含自动生成的单算子API执行接口头文件。

#include "aclnn_add_custom.h"
/home/HwHiAiUser/samples/operator/AddCustomSample/FrameworkLaunch/AddCustom/build_out/

在这里插入图片描述编译完之后开始安装
编译成功后打包为run文件
然后安装run,安装完整就可以用了
在这里插入图片描述
安装完就可以调用了

另一部分算子调用

/home/HwHiAiUser/samples/operator/AddCustomSample/FrameworkLaunch/AclNNInvocation/

请添加图片描述

├──input                                                 // 存放脚本生成的输入数据目录
├──output                                                // 存放算子运行输出数据和真值数据的目录
├── inc                           // 头文件目录 
│   ├── common.h                 // 声明公共方法类,用于读取二进制文件 
│   ├── operator_desc.h          // 算子描述声明文件,包含算子输入/输出,算子类型以及输入描述与输出描述 
│   ├── op_runner.h              // 算子运行相关信息声明文件,包含算子输入/输出个数,输入/输出大小等 
├── src 
│   ├── CMakeLists.txt    // 编译规则文件
│   ├── common.cpp         // 公共函数,读取二进制文件函数的实现文件
│   ├── main.cpp    // 单算子调用应用的入口
│   ├── operator_desc.cpp     // 构造算子的输入与输出描述 
│   ├── op_runner.cpp   // 单算子调用主体流程实现文件
├── scripts
│   ├── verify_result.py    // 真值对比文件
│   ├── gen_data.py    // 输入数据和真值数据生成脚本文件
│   ├── acl.json    // acl配置文件

运行结果
在这里插入图片描述

自定义算子编译部署后,会自动生成单算子API,可以直接在应用程序中调用。算子API的形式一般定义为“两段式接口”,形如:

aclnnStatus aclnnXxxGetWorkspaceSize(const aclTensor *src, ..., aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor);
aclnnStatus aclnnXxx(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream);

Xxx替换实际使用的函数

// 获取算子使用的workspace空间大小
aclnnStatus aclnnAddCustomGetWorkspaceSize(
const aclTensor *x, 
const aclTensor *y, 
const alcTensor *out, 
uint64_t workspaceSize, 
aclOpExecutor **executor);
// 执行算子
aclnnStatus aclnnAddCustom(void *workspace, 
int64_t workspaceSize, 
aclOpExecutor **executor, 
aclrtStream stream);

在这里插入图片描述

单算子API(aclnn)调用–关键代码

重要的就两个函数
在这里插入图片描述

bool OpRunner::RunOp()
{
    for (size_t i = 0; i < numInputs_; ++i) {
        auto size = GetInputSize(i);
        aclrtMemcpyKind kind = ACL_MEMCPY_HOST_TO_DEVICE;
        if (g_isDevice) {
            kind = ACL_MEMCPY_DEVICE_TO_DEVICE;
        }
        if (aclrtMemcpy(devInputs_[i], size, hostInputs_[i], size, kind) != ACL_SUCCESS) {
            ERROR_LOG("Copy input[%zu] failed", i);
            return false;
        }
        INFO_LOG("Copy input[%zu] success", i);
    }

    aclrtStream stream = nullptr;
    if (aclrtCreateStream(&stream) != ACL_SUCCESS) {
        ERROR_LOG("Create stream failed");
        return false;
    }
    INFO_LOG("Create stream success");

    size_t workspaceSize = 0;
    aclOpExecutor *handle = nullptr;
    auto ret =
        aclnnAddCustomGetWorkspaceSize(inputTensor_[0], inputTensor_[1], outputTensor_[0], &workspaceSize, &handle);
    if (ret != ACL_SUCCESS) {
        (void)aclrtDestroyStream(stream);
        ERROR_LOG("Get Operator Workspace failed. error code is %d", static_cast<int32_t>(ret));
        return false;
    }
    INFO_LOG("Execute aclnnAddCustomGetWorkspaceSize success, workspace size %lu", workspaceSize);

    if (workspaceSize != 0) {
        if (aclrtMalloc(&workspace_, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) {
            ERROR_LOG("Malloc device memory failed");
        }
    }

    ret = aclnnAddCustom(workspace_, workspaceSize, handle, stream);
    if (ret != ACL_SUCCESS) {
        (void)aclrtDestroyStream(stream);
        ERROR_LOG("Execute Operator failed. error code is %d", static_cast<int32_t>(ret));
        return false;
    }
    INFO_LOG("Execute aclnnAddCustom success");

    ret = aclrtSynchronizeStreamWithTimeout(stream, 5000);
    if (ret != SUCCESS) {
        ERROR_LOG("Synchronize stream failed. error code is %d", static_cast<int32_t>(ret));
        (void)aclrtDestroyStream(stream);
        return false;
    }
    INFO_LOG("Synchronize stream success");

    for (size_t i = 0; i < numOutputs_; ++i) {
        auto size = GetOutputSize(i);
        aclrtMemcpyKind kind = ACL_MEMCPY_DEVICE_TO_HOST;
        if (g_isDevice) {
            kind = ACL_MEMCPY_DEVICE_TO_DEVICE;
        }
        if (aclrtMemcpy(hostOutputs_[i], size, devOutputs_[i], size, kind) != ACL_SUCCESS) {
            INFO_LOG("Copy output[%zu] success", i);
            (void)aclrtDestroyStream(stream);
            return false;
        }
        INFO_LOG("Copy output[%zu] success", i);
    }

    (void)aclrtDestroyStream(stream);
    return true;
}
<think>好的,用户想在香橙设备上使用Ascend C进行算子开发。首先,我需要整理已有的引用内容,看看有没有相关的信息。引用[1]提到Ascend C对标CUDA,用于将Python模型编译成二进制文件,通过C/C++调用,这应该是Ascend C的主要用途。引用[2]和[4]涉及算子编译后的API调用和CANN路径,这些是开发过程中的具体步骤。引用[3]则说明香橙的软件支持,包括CANN框架,这对环境配置很重要。 接下来,用户需要教程或示例代码。根据引用内容,可能需要先配置CANN环境,确保香橙安装了必要的工具链。然后,编写Ascend C算子代码,编译生成二进制文件,最后调用测试。步骤可能需要包括环境准备、代码结构、编译命令、API调用示例。要注意引用中的路径和API使用方法,如引用[4]的路径和引用[2]的API调用部分。需要提醒用户检查环境是否安装正确,特别是CANN和Ascend Toolkit的位置。可能还需要涉及调试和性能优化的注意事项,但根据现有引用,暂时集中在基础流程上。最后生成相关问题,帮助用户深入理解相关开发步骤和可能遇到的问题。</think>### 香橙设备上使用Ascend C进行算子开发步骤 #### 1. **环境准备** - **安装CANN工具包** 根据引用[4],香橙AIpro的CANN工具包默认路径为`/usr/local/Ascend/ascend-toolkit/`。需确保已正确安装并配置环境变量: ```bash export ASCEND_TOOLKIT_HOME=/usr/local/Ascend/ascend-toolkit/latest export LD_LIBRARY_PATH=$ASCEND_TOOLKIT_HOME/lib64:$LD_LIBRARY_PATH ``` - **开发工具依赖** 安装必要的编译工具链(如gcc、cmake)和AI框架(如MindSpore),参考引用[3]中提到的开发工具支持。 --- #### 2. **Ascend C算子开发流程 ** ##### **步骤1:编写算子定义** 创建一个`.cpp`文件定义算子逻辑,例如实现向量加法: ```cpp // vec_add.cpp #include "acl/acl.h" extern "C" __global__ __aicore__ void VecAddKernel(float* x, float* y, float* z, int len) { for (int i = 0; i < len; ++i) { z[i] = x[i] + y[i]; } } ``` ##### **步骤2:编译算子为二进制** 使用CANN提供的编译器`ascendcc`编译算子: ```bash ascendcc -x cuda -o vec_add.o vec_add.cpp ``` --- #### 3. **调用算子** 根据引用[2],通过C语言API调用已编译的算子: ```cpp #include <stdio.h> #include "acl/acl.h" int main() { aclInit(nullptr); // 分配设备内存 float *x_dev, *y_dev, *z_dev; aclrtMalloc((void**)&x_dev, 1024 * sizeof(float)); aclrtMalloc((void**)&y_dev, 1024 * sizeof(float)); aclrtMalloc((void**)&z_dev, 1024 * sizeof(float)); // 调用算子 VecAddKernel<<<1, 1>>>(x_dev, y_dev, z_dev, 1024); // 同步并释放资源 aclrtSynchronizeStream(0); aclrtFree(x_dev); aclrtFree(y_dev); aclrtFree(z_dev); aclFinalize(); return 0; } ``` --- #### 4. **调试部署** - **日志性能分析** 使用`aclSetDumpPath`输出调试日志。 - **集成到AI模型** 通过MindSpore框架加载算子,参考引用[1]中模型编译调用的流程。 --- ###
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

二分掌柜的

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值