ATVOSS DeviceAdapter 详解host/device 桥接层的设计原理与 Run 运行流程【免费下载链接】atvossATVOSSAscend C Templates for Vector Operator Subroutines是一套基于Ascend C开发的Vector算子库致力于为昇腾硬件上的Vector类融合算子提供极简、高效、高性能、高拓展的编程方式。项目地址: https://gitcode.com/cann/atvossDeviceAdapter 是 ATVOSS 中位于 host 侧的 device 适配层类它把用户以表达式方式书写的算子配置KernelOp封装为可直接在昇腾硬件上运行的实体内部统一完成 ACL 资源管理、参数解析、Tiling 计算与 kernel 启动。本文以 DeviceAdapter.md 与 DeviceAdapter_Run.md 为骨架结合 device_adapter.h、tiling.h、arguments.h 等源码深入讲解其类设计、Run接口参数语义、底层五步执行流水线并给出可直接落地的完整代码示例帮助读者掌握 ATVOSS 算子从“表达式配置”到“device 上执行”的完整通路。一、DeviceAdapter 的功能定位ATVOSSAscend C Templates for Vector Operator Subroutines将 Vector 融合算子的开发抽象为三个层次Block 层由BlockBuilder承担把单个核的任务切分为多个 tile完成每次数据搬运与计算Kernel 层由KernelBuilder承担负责计算 Tiling 信息、依据 block ID 确定当前核需要处理的 GM 数据并下发给 BlockDevice 层由DeviceAdapter承担构造 device 适配层对象桥接 host 与 device。按照 device_adapter.h 中的类注释DeviceAdapter 是一个通用适配器generic adapter为不同的算子调用提供统一的 host 侧接口内部封装了 Acl 相关的资源管理并自动处理 kernel 调用。也就是说用户只需要面向表达式描述“算什么、怎么算”DeviceAdapter 负责把计算描述落到真实的昇腾执行环境中。其头文件为 device_adapter.h文档中给出的类原型如下template typename KernelOp class DeviceAdapter{ DeviceAdapter() {}; }类模板的唯一模板参数KernelOp是用户算子配置中kernel层的静态配置与调度策略即由Atvoss::Ele::KernelBuilderBlockOp实例化得到的类型。二、模板参数 KernelOp 说明参数名称参数类型输入/输出数据类型参数说明默认值KernelOp模板参数输入NAkernel 层的用户静态配置和调度策略NA从 device_adapter.h 可以看到KernelOp上还派生出一组关键类型别名它们决定了 DeviceAdapter 内部如何使用该配置using ExprMaker typename KernelOp::ScheduleClz::ExprMaker; using BlockOp typename KernelOp::ScheduleClz::BlockTemplate; using OpParam typename KernelOp::ScheduleCfgClz; template typename T using Tensor DeviceTensorT;ExprMaker表达式构造器用于在 host 侧重新实例化用户的计算表达式Compute()BlockOpkernel 层携带的 Block 模板类型供 Tiling 计算时逐级调用OpParamkernel/block 两级调度配置Tiling 信息的组合类型Tensordevice 侧的张量视图即 DeviceTensor它持有T* ptr_指针并提供GetPtr()、operator[]等访问方式IsTensor标记用于模板特化识别。在 examples/abs/abs.cpp 中可以看到真实的组装顺序先由BlockBuilder生成BlockOp再由KernelBuilder包装成KernelOp最终交给DeviceAdapterusing ArchTag Atvoss::Arch::DAV_3510; using BlockOp Atvoss::Ele::BlockBuilderAbsCompute, ArchTag, blockPolicy, Atvoss::Ele::DefaultBlockConfig; using KernelOp Atvoss::Ele::KernelBuilderBlockOp, kernelPolicy; using DeviceOp Atvoss::DeviceAdapterKernelOp;DeviceAdapter的无参构造函数DeviceAdapter() {}意味着该对象是轻量句柄通常作为局部对象实例化后直接调用Run。三、DeviceAdapter::Run 主运行接口3.1 函数原型template typename Args int64_t Run(const Args arguments, aclrtStream stream nullptr)Run是 device 适配层的主运行接口负责完成host 侧的参数解析与device 侧入参数据结构对象的准备随后触发 kernel 启动。3.2 参数说明参数名称参数类型输入/输出数据类型参数说明默认值Args模板参数输入NA用户的输入参数列表类型根据用户传入的参数实例化NAarguments函数形参输入Args用户传入的参数列表NAstream函数形参输入aclrtStream用户创建的 stream 流对象nullptr其中arguments通常由Atvoss::ArgumentsBuilder构建见 arguments.h它实际是一个二元组(inputOutput 元组, attrs 元组)。Run内部通过std::get0(arguments)取出输入输出参数元组继续处理。3.3 返回值说明返回值数据类型返回值说明int64_tdevice 层执行的结果0成功-1失败当 Tiling 计算失败时Run会打印[ERROR]: [Atvoss][Device] CalcParam failed!并返回 -1见 device_adapter.h。因此调用方应检查返回值以确认算子是否成功下发。四、Run 的源码级五步执行流水线Run的实现device_adapter.h可以分解为五个阶段下面结合源码逐步展开。4.1 表达式线性化还原用户计算语义auto expr ToLinearizerExpr(ExprMaker{}.template ComputeTensor()); using Expr typename decltype(expr)::Type; using Params Atvoss::Params_tExpr;这里用KernelOp携带的ExprMaker在 host 侧重新执行用户定义的Compute()得到计算表达式树后线性化linearize再通过Params_tExpr提取出全部PlaceHolder参数的静态描述参数序号、数据类型、ParamUsage流向等。这一步是后续参数准备与转换的“类型蓝图”。4.2 参数准备PrepareParamsauto argTuple std::get0(arguments); auto params PrepareParamsParams(argTuple);PrepareParams逐个把用户传入的实参按照PlaceHolder的序号构造为DeviceTensor或标量包装对象ConstructParam见 device_adapter.h当参数声明为TensorInputDtype而实参是Atvoss::Tensor时构造DeviceTensorT其内部直接引用用户传入的 device 指针ptr_ src.data()见 device_tensor.h当参数是标量如示例中的in3时直接构造标量类型。静态断言保证了Params的参数数量与用户实参元组大小严格一致不一致会在编译期报错见 device_adapter.h。4.3 动态参数计算TilingOpParam opParam; if (!CalculateTilingKernelOp(arguments, opParam)) { ... return -1; }CalculateTiling定义在 tiling.h它分两级调用调度策略先调用KernelOp::ScheduleClz::MakeScheduleConfig(arguments, cfg.kernelParam)计算 kernel 层配置如核数blockNum、每核处理单元数等参见 DefaultKernelConfig再调用BlockOp::ScheduleClz::MakeScheduleConfig(arguments, cfg.kernelParam, cfg.blockParam)计算 block 层配置如整 tile 数、尾 tile 元素数等参见 DefaultBlockConfig。arguments中的 attr如示例里的attr(dim, 5)在这里作为调度依据被消费。4.4 参数转换ConvertArgs 与 TransformArgsauto convertArgs ConvertArgsParams(params, argTuple);ConvertArgs依据Params中每个参数的位置与 usage把params与原始argTuple重新对齐成最终传给 kernel 的元组。随后LaunchKernelWithDataTuple会通过TransformArgs见 device_adapter.h完成最终形态转换标量类型原样转发std::forwardT(value)Tensor 类型取其 device 指针value.GetPtr()。TransformArgs有编译期static_assert只接受标量或IsTensor特化类型从类型系统上杜绝了非法入参。4.5 Kernel 启动LaunchKernelWithDataTupleLaunchKernelWithDataTupleKernelOp(opParam.kernelParam.blockNum, stream, opParam, convertArgs);启动入口是KernelCustomdevice_adapter.h它是一个__global__ __aicore__核函数声明了KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)即 AIV 核执行通过AscendC::Std::make_index_sequence展开参数元组并调用KernelWrapper最终执行KernelOp op; op.Run(cfg, args...)进入 kernel 层调度kernel/builder.h。此外源码中提供ATVOSS_DEBUG_MODE 2的性能剖析分支此时 kernel 会连续启动 200 次代码注释标注为 profiling run times便于观测算子耗时默认分支只启动一次。五、完整使用示例5.1 算子配置中的 DeviceAdapter以下示例完整来自 DeviceAdapter.md 与 DeviceAdapter_Run.md展示了一个 AddSubout in1 in2 - in3算子从 Compute 到 DeviceOp 的完整组装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; };要点说明PlaceHolderN, T, ParamUsage中的序号N决定参数在 kernel 实参列表中的位置ParamUsage::IN/OUT/IN_OUT声明数据流向参见 ParamUsage.mdArchTag Atvoss::Arch::DAV_3510指定目标昇腾芯片架构。5.2 host 侧调用 DeviceAdapter::Runtemplate typename InputDtype, typename OutputDtype static void Run() { /* ACL init and stream create */ ... 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; deviceOp.Run(arguments, stream); } int main(int argc, char const* argv[]) { Runfloat, float(); return 0; }几点实战说明Atvoss::Tensor的构造函数为Tensor(T* dataPtr, uint64_t* inputShape, size_t dims)shape 数组固定按 8 维描述未用到的维度补 0见 tensor.h示例中{{3, 4, 0, 0, 0, 0, 0, 0}}表示一个 3×4 的二维张量ArgumentsBuilder的inputOutput只接受Atvoss::Tensor与标量类型attr添加的键值对如dim会被 Tiling 计算消费不允许传入裸指针有static_assert保护见 arguments.hstream可省略默认nullptr但实际使用时建议传入通过aclrtCreateStream创建的 stream 流对象以便与其他 ACL 任务同步调度。5.3 与 ACL 初始化流程的完整衔接以 examples/abs/abs.cpp 为参照一次完整的DeviceAdapter::Run调用需要包裹在标准 ACL 生命周期中aclInit(nullptr)初始化 ACL结束时aclFinalize()aclrtSetDevice(deviceId)设置 deviceaclrtCreateContext(context, deviceId)创建 ContextaclrtCreateStream(stream)创建 StreamaclrtMalloc为输入/输出分配 device 内存ACL_MEM_MALLOC_HUGE_FIRSTaclrtMemcpy完成 Host→Device 数据拷贝用Atvoss::Tensor包装 device 指针并构建arguments实例化DeviceOp deviceOp;并调用deviceOp.Run(arguments, stream);aclrtSynchronizeStream(stream)同步再aclrtMemcpy拷回结果验证。该样例同时展示了返回值检查的实践Run返回 0 表示 kernel 成功下发结合 stream 同步即可安全读取输出。六、约束说明与易错点约束说明文档标注NA即无额外硬件/软件约束但使用时需保证Atvoss::Tensor包装的是已经通过aclrtMalloc分配的 device 内存而非 host 内存ArgumentsBuilder::inputOutput的实参类型必须为Atvoss::Tensor或标量禁止裸指针实参数量与Compute()中PlaceHolder的数量保持一致否则编译期static_assert会失败shape 维度数在 1~8 之间MAX_DIMS 8见 tensor.h。返回值检查务必检查Run的int64_t返回值0 成功、-1 失败Tiling 计算失败时返回 -1。性能剖析编译期定义ATVOSS_DEBUG_MODE 2会令Run连续启动 200 次 kernel 用于性能采样正式发布请勿开启。七、从文档到源码的进阶阅读路径类定义与 Run 实现device_adapter.hTiling 计算tiling.hDevice 侧张量视图device_tensor.hhost 侧参数构建arguments.hKernel 层构建器include/elewise/kernel/builder.hBlock 层构建器include/elewise/block/builder.h端到端样例examples/abs/abs.cpp、examples/muls/muls.cpp、examples/rms_norm/rms_norm.cpp相关 API 文档DeviceAdapter_Run.md、ArgumentsBuilder_inputOutput.md、ParamUsage.md通过本文可以完整理解用户在Compute()中书写表达式 →BlockBuilder/KernelBuilder逐级封装 →DeviceAdapter::Run完成表达式还原、参数准备、Tiling 计算、参数转换与 kernel 启动。掌握这条通路后即可基于 ATVOSS 快速开发并调试自己的 Vector 融合算子。【免费下载链接】atvossATVOSSAscend C Templates for Vector Operator Subroutines是一套基于Ascend C开发的Vector算子库致力于为昇腾硬件上的Vector类融合算子提供极简、高效、高性能、高拓展的编程方式。项目地址: https://gitcode.com/cann/atvoss创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考