ATVOSS DeviceAdapter::Run 深度解析Vector 算子 host 侧运行入口的完整调用链【免费下载链接】atvossATVOSSAscend C Templates for Vector Operator Subroutines是一套基于Ascend C开发的Vector算子库致力于为昇腾硬件上的Vector类融合算子提供极简、高效、高性能、高拓展的编程方式。项目地址: https://gitcode.com/cann/atvoss导读DeviceAdapter::Run是 CANN ATVOSSAscend C Templates for Vector Operator Subroutines设备适配层对外暴露的核心运行接口负责完成 host 侧参数解析、device 侧入参数据结构准备以及最终的内核Kernel启动。本文以 docs/api/DeviceAdapter_Run.md 为主干结合 include/elewise/device/device_adapter.h 及 tests/st/test_block_cast1.cpp 等源码与测试逐步拆解该接口的签名语义、内部执行流水线、前置参数封装方式与端到端调用姿势帮助开发者理解并正确使用 ATVOSS 的 host 侧算子运行入口。一、接口定位DeviceAdapter 在 ATVOSS 分层中的角色ATVOSS 的算子编程模型沿用了从底层向上逐层组装的思想用户先通过BlockBuilder描述块级计算BlockOp再经由KernelBuilder组装为内核级调度KernelOp而DeviceAdapterKernelOp则位于最上层负责桥接 host 与 device。构造DeviceAdapter详见 docs/api/DeviceAdapter.md仅产生一个适配层对象本身不涉及资源分配真正的工作发生在Run接口被调用时它完成 host 侧参数解析、tiling 计算含 workspace 规划、参数转换与 kernel 启动。从源码注释可以看到其设计定位include/elewise/device/device_adapter.hDeviceAdapter is a generic adapter that provides a host-side generic interface for different operator invocation. It encapsulates Acl-related resource management internally and automatically handles kernel invocation.即它面向不同算子提供统一的 host 侧泛型调用入口内部封装 ACL 相关资源管理并自动处理内核调用用户无需关心aclrtLaunch等底层细节。二、函数原型与签名语义Run是DeviceAdapter类模板的公开成员函数原型如下template typename KernelOp class DeviceAdapter { template typename Args int64_t Run(const Args arguments, aclrtStream stream nullptr); };类模板参数KernelOpkernel 层的用户静态配置与调度策略通常由Atvoss::Ele::KernelBuilderBlockOp实例化而来成员模板参数Args由用户传入的参数对象ArgumentsBuilder构建的结果自动推导无需显式指定stream用户创建的 ACL stream 流对象默认为nullptr由 KernelCustom 内部处理此时依赖默认流语义。参数说明参数名称参数类型输入/输出数据类型参数说明默认值Args模板参数输入NA用户的输入参数列表类型根据用户传入的参数实例化NAarguments函数形参输入Args用户传入的参数列表由ArgumentsBuilder构建NAstream函数形参输入aclrtStream用户创建的 stream 流对象nullptr返回值说明返回值数据类型返回值说明int64_tdevice 层执行的结果0成功-1失败需要特别注意的是Run返回的-1特指调度配置阶段失败。从 include/elewise/device/device_adapter.h 可以看到当CalculateTilingKernelOp(arguments, opParam)返回false时Run会打印错误并立即返回-1kernel 启动本身是异步的真正的计算错误需要通过后续aclrtSynchronizeStream等同步手段捕获。三、Run 的内部执行流水线从参数到 Kernel 启动Run的实现include/elewise/device/device_adapter.h内部由四个阶段组成每一步都对应独立的模板元编程辅助函数template typename Args int64_t Run(const Args arguments, aclrtStream stream nullptr) { auto expr ToLinearizerExpr(ExprMaker{}.template ComputeTensor()); using Expr typename decltype(expr)::Type; using Params Atvoss::Params_tExpr; auto argTuple std::get0(arguments); // 1. prepare Param auto params PrepareParamsParams(argTuple); // 2. calc dynamic param tiling / workspace OpParam opParam; if (!CalculateTilingKernelOp(arguments, opParam)) { printf([ERROR]: [Atvoss][Device] CalcParam failed!\n); return -1; } // 3. kernel launch auto convertArgs ConvertArgsParams(params, argTuple); LaunchKernelWithDataTupleKernelOp(opParam.kernelParam.blockNum, stream, opParam, convertArgs); return 0; }3.1 阶段一表达式重建与参数类型推导auto expr ToLinearizerExpr(ExprMaker{}.template ComputeTensor()); using Expr typename decltype(expr)::Type; using Params Atvoss::Params_tExpr;ExprMaker即KernelOp::ScheduleClz::ExprMaker见 include/elewise/device/device_adapter.h它重新以 device 侧Tensor模板DeviceTensorT执行用户写的Compute()表达式将其线性化为可编译的表达式类型Expr随后通过Atvoss::Params_tExpr定义于 include/expression/expr_template.h在编译期提取出该表达式涉及的全部占位符参数PlaceHolder列表。这意味着Run不需要用户手动声明参数个数与顺序参数结构完全由表达式推导而来这正是 ATVOSS 极简编程方式的体现。3.2 阶段二PrepareParams 组装 device 侧参数对象auto argTuple std::get0(arguments); auto params PrepareParamsParams(argTuple);arguments是ArgumentsBuilder::build()的结果本质是一个二元组std::make_tuple(inputOutput, attrs)见 include/utils/arguments/arguments.h因此std::get0(arguments)取出的是用户传入的输入输出参数元组。PrepareParams依据编译期推导出的Params类型列表按ParamType::number - 1从argTuple中取出对应位置的实参构造出统一的DeviceTensorT或标量参数对象见ConstructParaminclude/elewise/device/device_adapter.h若占位符声明为标量类型、而实参是Atvoss::Tensor则用实参数据构造标量Tensor否则直接以实参类型实例化参数对象。3.3 阶段三CalculateTiling 计算动态调度参数OpParam opParam; if (!CalculateTilingKernelOp(arguments, opParam)) { return -1; }CalculateTiling定义于 include/elewise/device/tiling.h内部依次调用两级调度配置Kernel 级KernelOp::ScheduleClz::MakeScheduleConfig(args, cfg.kernelParam)产出 block 数量kernelParam.blockNum等内核级调度信息Block 级BlockOp::ScheduleClz::MakeScheduleConfig(args, cfg.kernelParam, cfg.blockParam)产出 block 内的 tiling 与 workspace 规划。两级任一失败都会打印[ERROR]: [Atvoss][Device] ...并返回falseRun随即返回-1。关于这两级配置接口的详细语义可参考 docs/api/BaseKernelSchedule_MakeScheduleConfig.md 与 docs/api/BaseBlockSchedule_MakeScheduleConfig.md。3.4 阶段四ConvertArgs 与 Kernel 启动auto convertArgs ConvertArgsParams(params, argTuple); LaunchKernelWithDataTupleKernelOp(opParam.kernelParam.blockNum, stream, opParam, convertArgs);ConvertArgs依据Params中CheckVarNumIndex 1的位置信息将参数元组重排为 device 侧 kernel 期望的实参序列随后LaunchKernelWithDataTupleinclude/elewise/device/device_adapter.h通过std::applyTransformArgs完成最后一层转换标量参数原样转发Tensor 参数调用DeviceTensor::GetPtr()退化为裸设备指针T*。TransformArgs通过static_assert约束入参只能是标量类型或Atvoss::Tensor特化include/elewise/device/device_adapter.h。最终内核入口为KernelCustominclude/elewise/device/device_adapter.h它被声明为__global__ __aicore__并使用KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)指定 AI Vector 任务类型通过三元组 launch 语法KernelCustomKernelOp, OpParamblockNum, nullptr, stream(cfg, transformedArgs)提交到指定 streamKernelWrapper则在 AI Core 上实例化KernelOp并解包参数元组调用op.Run(cfg, args...)。另外源码中保留了 profiling 开关当编译期宏ATVOSS_DEBUG_MODE 2时Run会连续启动 200 次 kernel 以便性能采样见 include/elewise/device/device_adapter.h常规构建下只启动一次。四、调用前置ArgumentsBuilder 参数封装Run的第一个入参arguments必须由Atvoss::ArgumentsBuilder构建。构建器include/utils/arguments/arguments.h提供链式 APIauto arguments Atvoss::ArgumentsBuilder{} .inputOutput(in1, in2, in3, out) // 输入输出参数 .attr(dim, 5) // 算子属性 .build(); // 产出 std::tupleinputOutput, attrs其中inputOutput环节带有两层编译期约束static_assert参数不允许是指针类型Pointer types are not allowed in inputOutput parameters参数只能是Atvoss::Tensor特化或标量类型Only Atvoss::Tensor and scalar types are allowed in inputOutput parameters。attr将键值对封装为AttrMapKey, Value通过MakeAttr可以多次链式追加build()返回std::make_tuple(inOutCollector.inputOutput, attrCollector.attrs)与Run内部std::get0(arguments)的取用方式一一对应。五、完整使用示例AddSub 三输入算子以下示例完整继承自 docs/api/DeviceAdapter_Run.md 并补充了关键注释。它演示了一个out in1 in2 - in3的融合算子从配置声明到Run调用的全过程template typename InputDtype, typename OutputDtype struct AddSubConfig { struct AddSubCompute { template template typename class Tensor __host_aicore__ constexpr auto Compute() const { auto in1 Atvoss::PlaceHolder1, TensorInputDtype, Atvoss::ParamUsage::IN(); auto in2 Atvoss::PlaceHolder2, TensorInputDtype, Atvoss::ParamUsage::IN(); auto in3 Atvoss::PlaceHolder3, InputDtype, Atvoss::ParamUsage::IN(); // 标量输入 auto out Atvoss::PlaceHolder4, TensorOutputDtype, Atvoss::ParamUsage::OUT(); return (out in1 in2 - in3); }; }; using ArchTag Atvoss::Arch::DAV_3510; using BlockOp Atvoss::Ele::BlockBuilderAddSubCompute, ArchTag; using KernelOp Atvoss::Ele::KernelBuilderBlockOp; using DeviceOp Atvoss::DeviceAdapterKernelOp; // 组装 DeviceAdapter }; template typename InputDtype, typename OutputDtype static void Run() { /* ACL init and stream create */ ... // 构造 device 侧 Tensor裸指针 8 维 shape 数组 实际维数 Atvoss::TensorInputDtype in1(deviceIn1, {{3, 4, 0, 0, 0, 0, 0, 0}}, 2); Atvoss::TensorInputDtype in2(deviceIn2, {{3, 4, 0, 0, 0, 0, 0, 0}}, 2); InputDtype in3 5.0; // 标量参数 Atvoss::TensorOutputDtype out(deviceOut, {{3, 4, 0, 0, 0, 0, 0, 0}}, 2); // 构建参数对象 auto arguments Atvoss::ArgumentsBuilder{}.inputOutput(in1, in2, in3, out).attr(dim, 5).build(); using DeviceOp typename AddSubConfigInputDtype, OutputDtype::DeviceOp; DeviceOp deviceOp; // 构造适配层对象 // 核心调用Run(arguments, stream) deviceOp.Run(arguments, stream); } int main(int argc, char const* argv[]) { Runfloat, float(); return 0; }要点说明PlaceHolder的第一个模板参数是全局唯一编号Run内部通过该编号与ArgumentsBuilder中参数的位置顺序对齐ParamType::number - 1索引ParamUsage::IN / OUT / IN_OUT决定了参数在调度阶段的 CopyIn / CopyOut 归类见 include/elewise/device/device_adapter.h标量输入如in3同样通过inputOutput传入不需要单独接口stream可省略默认nullptr但生产代码建议显式传入用户创建的 stream 以保证与上层调用流同步。六、端到端实战测试用例中的标准调用环境tests/st目录下的算子测试如 tests/st/test_block_cast1.cpp给出了Run之外的完整 host 侧环境准备流程可作为实战模板ACL 初始化aclInit(nullptr)并在退出时aclFinalize()设置 deviceaclrtSetDevice(deviceId)结束调用aclrtResetDevice创建 contextaclrtCreateContext(context, deviceId)创建 streamaclrtCreateStream(stream)分配设备内存aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)将裸指针传入Atvoss::TensorTHost → Device 拷贝aclrtMemcpy(..., ACL_MEMCPY_HOST_TO_DEVICE)构造参数并调用Runuint64_t shapeArray[MAX_DIM] {0}; std::copy(shape.begin(), shape.end(), shapeArray); Atvoss::TensorT1 t1(deviceInput, shapeArray, shape.size()); Atvoss::TensorT2 t2(deviceOutput, shapeArray, shape.size()); auto arguments Atvoss::ArgumentsBuilder{}.inputOutput(t1, t2).build(); using DeviceOp typename CastConfigT1, T2, tileShapeLen::DeviceOp; DeviceOp deviceOp; deviceOp.Run(arguments, stream);同步并取回结果aclrtSynchronizeStream(stream)之后aclrtMemcpy(..., ACL_MEMCPY_DEVICE_TO_HOST)。其中第 7 步就是Run接口的标准调用现场测试通过CastConfig将DeviceOp定义为Atvoss::DeviceAdapterKernelOp与第五节示例中的用法完全一致可在 tests/st/test_block_cast1.cpp 等十余个 cast 用例中反复验证。七、约束与注意事项参数类型约束ArgumentsBuilder::inputOutput只接受Atvoss::Tensor与标量类型禁止裸指针TransformArgs在 kernel 启动前会再次校验同样约束参数顺序约束PlaceHolderN, ...的编号必须与inputOutput(...)中的实参顺序自 1 起连续对应否则ConstructParam会取错参数或触发编译期断言返回值语义0表示调度配置成功并已提交 kernel-1表示 tiling 计算失败kernel 异步执行的实际结果需配合aclrtSynchronizeStream确认错误定位源码中所有失败分支均会打印带[ERROR]: [Atvoss][Device]前缀的错误信息可通过日志快速定位是 kernel 级还是 block 级MakeScheduleConfig失败stream 生命周期Run只消费 stream不负责创建与销毁stream 的创建/销毁由调用方如测试中的ReleaseSource守卫管理。八、小结DeviceAdapter::Run是 ATVOSS host 侧唯一对外暴露的运行入口它将“表达式重建 → 参数类型推导 → device 参数准备 → 两级 tiling 计算 → 参数转换 → 异步 kernel 启动”这一整条链路收敛为一次模板化调用。开发者只需沿用“BlockBuilder → KernelBuilder → DeviceAdapter”的组装模式配合ArgumentsBuilder封装参数即可在数十行代码内完成一个 Vector 融合算子的 host 侧运行而Run内部的元编程与 ACL 细节则由 ATVOSS 统一接管。进一步阅读DeviceAdapter 构造接口、ArgumentsBuilder 构建说明、ArgumentsBuilder::build、KernelBuilder、BlockBuilder以及端到端样例 examples 与 tests/st 中的完整测试工程。【免费下载链接】atvossATVOSSAscend C Templates for Vector Operator Subroutines是一套基于Ascend C开发的Vector算子库致力于为昇腾硬件上的Vector类融合算子提供极简、高效、高性能、高拓展的编程方式。项目地址: https://gitcode.com/cann/atvoss创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考