企业官网建设流程全解析

CANN ops-math ClipByValue 算子深度解析：功能、参数、GE IR 图模式调用与源码实现

【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math

ClipByValue（裁剪取值）是 CANN ops-math 数学算子库中用于将张量元素限制到[clip_value_min, clip_value_max]区间的逐元素算子，其计算逻辑与 TensorFlow 框架的ClipByValue算子完全对齐。本文以 conversion/clip_by_value/README.md 为骨架，结合仓库内算子原型、Shape 推导、Tiling 与 Kernel 源码及测试用例，完整讲解该算子的支持产品范围、参数语义、GE IR 图模式调用方式，并深入剖析其从算子注册到 NPU 上执行的完整实现链路，帮助读者在 CANN 环境中快速上手并理解底层原理。

一、功能说明与计算公式

ClipByValue 算子的核心功能是：将输入张量x的所有元素限制在[clip_value_min, clip_value_max]区间内。具体规则为：

• 若元素x_i > clip_value_max，则输出被限制为clip_value_max；
• 若元素x_i < clip_value_min，则输出被限制为clip_value_min；
• 否则，输出等于元素本身。

其计算公式如下：

$$ {y}{i} = max(min({{x}{i}},{max_value}{i}),{min_value}{i}) $$

即对每个元素先与上界clip_value_max取min，再与下界clip_value_min取max。该语义在仓库的 golden 参考实现中得到了完全一致的印证，tests/assets/golden.py 中clip_by_value_golden函数的实现为：

min_ = np.minimum(x, clip_value_max) res = np.maximum(min_, clip_value_min)

从算子原型看，该算子兼容 TensorFlow 框架的同名算子 ClipByValue（见 op_graph/clip_by_value_proto.h 中注释 "Compatible with the TensorFlow operator ClipByValue"），在 CANN 场景下常用于梯度裁剪（gradient clipping）、数值稳定性保护等需要对数据范围做硬性约束的环节，防止训练过程中梯度爆炸或中间结果溢出。

二、产品支持情况

原文档给出了 ClipByValue 算子在 CANN 各产品形态上的支持矩阵，归纳如下：

值得注意的是，Atlas 200I/500 A2 推理产品当前不支持该算子，在实际部署前需要确认目标硬件平台是否在支持列表中。

这一支持情况与仓库源码中的算子配置是呼应的：op_host/clip_by_value_def.cpp 中通过this->AICore().AddConfig("ascend950", aicoreConfig)与this->AICore().AddConfig("ascend350", aicoreConfig)分别注册了 ascend950 与 ascend350 两个平台的 AICore 配置，对应产品矩阵中的 Ascend 950 系列与 Atlas A2/A3 系列（arch35），而 op_kernel 与 tiling 的实现均位于arch35目录下，从源码结构上印证了算子当前主要面向这两类架构提供 AscendC 内核实现。

三、参数说明

ClipByValue 算子共包含 3 个输入和 1 个输出，参数语义如下表所示：

3.1 原型层面支持的数据类型

从算子原型注册 op_graph/clip_by_value_proto.h 看，REG_OP(ClipByValue)中为x、clip_value_min、clip_value_max与输出y注册了完全一致的数据类型集合，包括：

DT_COMPLEX128, DT_COMPLEX64, DT_DOUBLE, DT_FLOAT, DT_FLOAT16, DT_INT16, DT_INT32, DT_INT64, DT_INT8, DT_QINT32, DT_QINT8, DT_QUINT8, DT_UINT16, DT_UINT8, DT_BF16, DT_COMPLEX32

即在 GE（Graph Engine）构图层面，该算子支持从 INT8 到 COMPLEX128 的 16 种数据类型，且三个输入与输出必须保持相同 dtype。同时原型注释明确说明：当输入为 bfloat16、float16、float32、int32 或 int64 时，支持广播（broadcasting）操作，即clip_value_min、clip_value_max的 shape 可以与x不同，通过广播机制逐元素对齐。

3.2 实际 AICore 内核支持的数据类型

需要特别说明的是，算子原型注册的完整类型集合与最终 NPU 内核实际支持的类型集合并不完全一致。在 op_host/clip_by_value_def.cpp 的 OpDef 注册中，四个参数的数据类型被限定为：

ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64

对应 op_host/arch35/clip_by_value_tiling.cpp 中DoOpTiling()的分支判断：仅对DT_FLOAT16、DT_BF16、DT_FLOAT、DT_INT32、DT_INT64五种 dtype 生成 tiling，其余类型会通过OP_LOGE_FOR_INVALID_DTYPE报错并返回GRAPH_FAILED。这与 op_host/config/ascend950/clip_by_value_binary.json（以及 ascend350 同名配置文件）中预编译二进制清单一致——该清单仅为 bfloat16、float16、float32、int32、int64 五种 dtype 各生成一个bin_filename条目。

因此在实际使用中，可直接在 NPU 上高效执行的数据类型为 FLOAT16、BFLOAT16、FLOAT、INT32、INT64 五种，数据格式统一为 ND（format_match_mode为 FormatAgnostic）；其他类型在构图层面合法，但取决于目标硬件的实现选择。这是 README 参数表之外，从源码中可以得到的重要补充信息。

3.3 输入输出约束

• 三个输入（x、clip_value_min、clip_value_max）与输出y的数据类型必须保持一致。Tiling 源码中的CheckDtype()方法（clip_by_value_tiling.cpp）会显式校验四者 dtype 完全一致，否则报错 "The dtype of x, clip_value_min, clip_value_max and y must be the same" 并返回失败。
• 输入输出的数据格式均为 ND。
• clip_value_min与clip_value_max的 shape 需与x满足广播规则（详见下节 Shape 推导）。

四、约束说明

原文档声明该算子"无约束"，即算子本身对输入 shape、rank 等没有额外的静态限制。这一点与 OpDef 中的 AICore 配置一致：clip_by_value_def.cpp 中设置了DynamicRankSupportFlag(true)（支持动态 rank）与DynamicShapeSupportFlag(true)（支持动态 shape），DynamicCompileStaticFlag(true)表示动态编译静态标志，PrecisionReduceFlag(true)表示允许精度降级优化。

不过，"无约束"是相对算子自身语义而言的，实际运行仍受以下通用规则约束：

1. 广播合法性：clip_value_min/clip_value_max与x的 shape 必须满足 NumPy 风格的广播规则，否则 Shape 推导阶段会直接失败（见下文测试用例中的GRAPH_FAILED场景）。
2. 平台支持约束：算子需在第二节产品支持矩阵所列的平台上运行，Atlas 200I/500 A2 推理产品不支持。
3. 数值语义约束：当clip_value_min > clip_value_max时，由于计算顺序是"先取 min 再取 max"，实际结果将由clip_value_min主导，这是由公式y_i = max(min(x_i, max_value_i), min_value_i)决定的数学性质，使用时需自行保证上下界语义正确。

五、调用说明：GE IR 图模式调用

原文档给出的调用方式为图模式调用，即通过算子 IR 构图的方式调用 ClipByValue 算子。对应的完整可运行样例位于 examples/test_geir_clip_by_value.cpp，其使用的算子 IR 定义见 op_graph/clip_by_value_proto.h。

注：原 README 中的相对链接./examples/test_geir_clip_by_value.cpp与./op_graph/clip_by_value_proto.h分别对应仓库根目录下的 conversion/clip_by_value/examples/test_geir_clip_by_value.cpp 与 conversion/clip_by_value/op_graph/clip_by_value_proto.h。

5.1 算子 IR 定义

clip_by_value_proto.h 通过REG_OP宏在ge命名空间中注册算子，关键定义如下：

REG_OP(ClipByValue) .INPUT(x, TensorType({DT_COMPLEX128, ..., DT_COMPLEX32})) .INPUT(clip_value_min, TensorType({DT_COMPLEX128, ..., DT_COMPLEX32})) .INPUT(clip_value_max, TensorType({DT_COMPLEX128, ..., DT_COMPLEX32})) .OUTPUT(y, TensorType({DT_COMPLEX128, ..., DT_COMPLEX32})) .OP_END_FACTORY_REG(ClipByValue)

5.2 图模式调用示例代码解析

样例程序 test_geir_clip_by_value.cpp 的完整执行流程如下：

第 1 步：创建 Graph 并构造算子节点。在CreateOppInGraph()函数中，通过op::ClipByValue("clipByValue1")创建算子实例，输入 shape 定义为xShape = {(1, 1, 1, 1)}，clip_value_min与clip_value_max的 shape 均为{1}（标量上下界，广播作用于 4 维输入）：

auto clipByValue1 = op::ClipByValue("clipByValue1"); std::vector<int64_t> xShape = {(1, 1, 1, 1)}; ADD_INPUT(1, x, inDtype, xShape); ADD_INPUT(2, clip_value_min, inDtype, {1}); ADD_INPUT(3, clip_value_max, inDtype, {1}); outputs.push_back(clipByValue1);

其中ADD_INPUT宏的核心动作包括：创建op::Data占位节点、构造TensorDesc（FORMAT_ND+ 指定 dtype）、将占位节点接入算子输入（set_input_x等）、生成全 1 填充的输入 Tensor 并加入 graph。示例使用DT_FLOAT作为输入 dtype（DataType inDtype = DT_FLOAT）。

第 2 步：初始化 GE 全局环境。通过ge::GEInitialize传入全局选项（示例中指定ge.exec.deviceId=0与ge.graphRunMode=1）：

std::map<AscendString, AscendString> global_options = { {"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"} }; Status ret = ge::GEInitialize(global_options);

第 3 步：创建 Session 并将计算图加入会话。使用ge::Session创建推理会话，调用session->AddGraph(graph_id, graph, graph_options)将 ClipByValue 计算图注册到会话中，并可通过aclgrphDumpGraph(graph, "./dump", ...)将构图结果导出为 txt 文件用于调试。

第 4 步：运行图并获取结果。调用session->RunGraph(graph_id, input, output)执行算子，运行成功后SaveInputOutput()会将输入输出张量分别写入二进制文件（tc_ge_irrun_test_0008_npu_input_*.bin/tc_ge_irrun_test_0008_npu_output_*.bin），并将输出结果以浮点形式逐元素打印到控制台：

LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);

第 5 步：收尾清理。读取 GE 的 error/warning 消息（GEGetErrorMsgV2/GEGetWarningMsgV2），调用ge::GEFinalize()结束 GE 会话。

该样例完整演示了"算子 IR 定义 → 图构建 → 会话执行 → 结果落盘"的 GE IR 单算子调用范式，是理解 CANN 图模式下自定义算子调用流程的绝佳入门材料。

5.3 框架侧接入

除 GE IR 直接构图外，仓库还提供了框架插件用于将第三方框架的 ClipByValue 算子映射到本实现：

• framework/clip_by_value_tf_plugin.cpp：注册 TensorFlow 框架的ClipByValue算子（FrameworkType(TENSORFLOW)、OriginOpType("ClipByValue")），采用AutoMappingByOpFn自动参数映射，ImplyType(ImplyType::TVM)。
• framework/clip_onnx_plugin.cpp：注册 ONNX 框架侧的算子映射插件。

这从源码层面印证了算子与 TensorFlow 生态的兼容性定位，也表明在昇腾推理/训练场景中，框架侧模型经过解析后即可落图到该算子实现。

六、源码级实现剖析：从 Shape 推导到 NPU Kernel

为了更深入理解算子行为，下面按执行链路的五个环节逐一剖析其源码实现。整个算子位于仓库 conversion/clip_by_value 目录下，CMakeLists.txt统一组织其编译。

6.1 算子定义（OpDef）

op_host/clip_by_value_def.cpp 通过OP_ADD(ClipByValue)注册 OpDef，核心信息：

• 三输入一输出均为REQUIRED（必选参数）；
• 支持 dtype：FLOAT16、BF16、FLOAT、INT32、INT64；
• 支持 format：ND；
• AICore 配置：动态编译静态标志开启、动态 rank/shape 支持开启、无需额外 support 检查、允许精度降级，并注册到 ascend950 与 ascend350 两个平台，扩展配置opFile.value = "clip_by_value_apt"。

6.2 Shape 推导（InferShape）

op_host/clip_by_value_infershape.cpp 的推导逻辑极为简洁：

IMPL_OP_INFERSHAPE(ClipByValue).InferShape(Ops::Base::InferShape4Broadcast);

即复用基础库的广播式 Shape 推导InferShape4Broadcast：输出y的 shape 等于三个输入按广播规则对齐后的结果 shape。这正是算子支持clip_value_min/clip_value_max与xshape 不同（广播）的根本机制。

该行为在单元测试 tests/ut/op_host/test_clip_by_value_infershape.cpp 中得到充分验证，典型用例包括：

测试同时覆盖了 ND 与 NHWC 格式下的动态 shape（-1）、未知 rank（-2）、零维（0）等多种边界情况，说明 Shape 推导对动态形状场景具备良好的鲁棒性。

6.3 Tiling 计算（算力调度决策）

op_host/arch35/clip_by_value_tiling.cpp 负责在 Host 侧完成算子切分策略（Tiling）决策：

• dtype 校验：CheckDtype()强制要求x、clip_value_min、clip_value_max、y四者 dtype 完全一致；
• 按 dtype 分发：FLOAT16/BF16 使用ClipByValueCompute<half>::OpDag，FLOAT 使用ClipByValueCompute<float>::OpDag，INT32/INT64 分别使用int32_t/int64_t特化，统一通过模板类BroadcastBaseTiling<...>完成带广播语义的 Tiling，并依据GetSchMode()生成对应的tilingKey；
• 编译期平台信息：TilingPrepareForClipByValue在编译期获取 AIV 核数（GetCoreNumAiv）与 UB 内存大小（GetCoreMemSize(CoreMemType::UB)），写入ClipByValueCompileInfo{coreNum, ubSize}（定义见 arch35/clip_by_value_tiling.h）；
• 注册方式：通过IMPL_OP_OPTILING(ClipByValue)注册 Tiling 入口，并用REGISTER_OPS_TILING_TEMPLATE(ClipByValue, ClipByValueTiling, 0)将ClipByValueTiling以优先级 0 挂入数学算子的通用 Tiling 模板注册表。

由于底层复用了atvoss/broadcast/broadcast_tiling.h的BroadcastBaseTiling，算子天然获得了对广播输入与多核并行的切分能力，无需针对每个 shape 单独编写切分逻辑。

6.4 Kernel 实现（Device 侧计算）

NPU 侧的 AscendC 内核位于 op_kernel/arch35/clip_by_value.cpp，核心入口是一个以schMode为模板参数的__global__ __aicore__函数，整体结构为"DAG 调度 + 广播切分"模式：

template <uint64_t schMode> __global__ __aicore__ void clip_by_value( GM_ADDR x, GM_ADDR clipValueMin, GM_ADDR clipValueMax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { using OpDag = ClipByValueOp::ClipByValueCompute<DTYPE_X>::OpDag; BroadcastSch<schMode, OpDag> sch(tiling); sch.Process(x, clipValueMin, clipValueMax, y); }

算子实际计算逻辑由 op_kernel/arch35/clip_by_value_dag.h 中的模板描述：

template <typename T> struct ClipByValueCompute { using OpInputX = Bind<Vec::CopyInBrc<T>, Placeholder::In0<T>>; // 广播搬入 x using OpInputMin = Bind<Vec::CopyInBrc<T>, Placeholder::In1<T>>; // 广播搬入 clip_value_min using OpInputMax = Bind<Vec::CopyInBrc<T>, Placeholder::In2<T>>; // 广播搬入 clip_value_max using OpClipRes = Bind<ClipByValueFused<T>, OpInputX, OpInputMin, OpInputMax>; // 融合计算 using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpClipRes>; // 写回 y using Outputs = Elems<OpCopyOut>; using OpDag = DAGSch<Outputs>; };

计算核心ClipByValueFused<T>在向量寄存器级别按max(min(x, max), min)语义完成元素操作：先用Vec::Min将x与上界取小，再用Vec::Max与下界取大。代码中针对int64_t类型还提供了基于VECTOR_REG_WIDTH_2XVL的特化路径（见 clip_by_value_dag.h），确保不同位宽的数据都能以合适的向量长度高效计算。整个过程通过BroadcastSch调度器结合 Tiling 结果，在多个 AIV 核上并行处理广播展开后的数据块。

6.5 预编译二进制与构建配置

针对 ascend950 与 ascend350 平台，仓库分别提供了预编译二进制清单：

• op_host/config/ascend950/clip_by_value_binary.json
• op_host/config/ascend350/clip_by_value_binary.json

每个清单包含 5 个op_list条目，分别对应 bfloat16、float16、float32、int32、int64 五种 dtype；每个条目标注了bin_filename（编译产物指纹文件名）、输入输出 name/index/dtype/format（均为 ND、format_match_mode: FormatAgnostic）以及 shape 模板（[-2]，表示任意动态 shape）。该文件是算子二进制产物与 shape/dtype 模板的映射索引，用于运行时的算子二进制匹配与加载。

七、测试与验证体系

7.1 单元测试（UT）

• Shape 推导测试：tests/ut/op_host/test_clip_by_value_infershape.cpp 覆盖上文表格所列的广播、动态 shape、未知 rank、标量输入等 20 余个用例，通过gtest+InfershapeContextPara构造算子上下文并断言输出 shape 与返回码。
• Tiling 测试：tests/ut/op_host/arch35/test_clip_by_value_tiling.cpp 用于验证 arch35 平台上的 Tiling 结果正确性。

7.2 黄金数据与 ST 测试

• Golden 参考实现：tests/assets/golden.py 以numpy.minimum/numpy.maximum实现了算子参考逻辑，作为结果比对的基准；对 bfloat16 输入会先提升到 float32 计算再转回 bfloat16，保证低精度场景下的数值精度。
• ST（系统级）测试：tests/st/arch35/ttk_kernel_clip_by_value_st.csv 定义了算子内核在 arch35 上的系统级测试用例，与 golden 脚本配合完成端到端正确性验证。

八、小结

ClipByValue 是 CANN ops-math 中一个语义简洁但工程链路完整的逐元素裁剪算子：

• 使用层面：掌握公式y_i = max(min(x_i, max_value_i), min_value_i)、三输入一输出的参数语义、ND 格式与五种内核 dtype（FLOAT16/BF16/FLOAT/INT32/INT64），并通过 test_geir_clip_by_value.cpp 复现 GE IR 图模式调用全流程；
• 实现层面：沿 clip_by_value_proto.h（原型）→ clip_by_value_def.cpp（OpDef）→ clip_by_value_infershape.cpp（广播 Shape 推导）→ clip_by_value_tiling.cpp（Tiling）→ clip_by_value.cpp（AscendC Kernel）的调用链，可以完整理解一个昇腾算子从注册、推导、切分到上核执行的标准化流程；
• 验证层面：利用 test_clip_by_value_infershape.cpp 与 golden.py 可快速回归算子语义与边界行为。

对于需要在昇腾 NPU 上执行梯度裁剪、数据范围约束等任务的开发者，本算子开箱即用；对于希望深入 CANN 算子开发流程的读者，ClipByValue 则是一份结构清晰、可完整对照源码研读的优质参考实现。

【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math