企业官网建设流程全解析

CANN ops-math TransposeV2 算子实战指南：维度置换接口、Tiling 实现与 aclnn 调用详解

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

本文基于 CANN ops-math 仓库中conversion/transpose_v2模块的官方 README 及配套源码展开，系统讲解 TransposeV2 算子的功能语义、参数与数据类型约束、通过aclnnPermute两段式接口调用的完整实操流程，并结合 host 侧 Tiling 分核分块策略与 device 侧 kernel 分发逻辑，给出该算子“参数校验—Tiling 计算—kernel 执行”全链路的源码级解读，帮助读者在 NPU 上正确集成并深入理解维度置换算子的实现原理。

1. 算子概述与产品支持情况

TransposeV2 是 CANN 数学算子库（ops-math）conversion目录下的一个 AICore 算子，用于实现张量的维度置换（Permutation）操作：按照指定顺序重新排列输入张量的维度。官方文档位于 conversion/transpose_v2/README.md。

1.1 产品支持矩阵

按照 README 中的产品支持情况表：

从源码结构看，上述支持列表与算子注册配置一一对应。在 transpose_v2_def.cpp 中，算子通过OpDef注册了四类 AICore 配置：

• ascend910b（Atlas A2 训练/推理系列）
• ascend910_93（Atlas A3 训练/推理系列）
• kirinx90与kirin9030（Kirin X90/Kirin 9030 处理器系列）

其中 Kirin 平台使用单独的配置函数GetKirinCoreConfig()（见 transpose_v2_def.cpp），额外开启了动态编译、动态 rank、动态 shape 等标志，并且数据类型只保留FLOAT16与FLOAT，不包含BFLOAT16——这与 README 中“Kirin X90/Kirin 9030 处理器系列产品：数据类型支持 FLOAT、FLOAT16”的说明完全吻合。

各平台对应的算子配置目录也直接体现了硬件目标，位于 op_host/config 下：ascend910b/、ascend910_93/、kirinx90/、kirin9030/四个子目录，各自包含transpose_v2_binary.json（算子属性配置）与transpose_v2_simplified_key.ini（Tiling key 简化配置）。

2. 功能说明：维度置换的语义

README 给出的功能定义是：

实现张量的维度置换（Permutation）操作，按照指定的顺序重新排列输入张量的维度。如输入self是 shape 为 [2, 3, 5] 的 tensor，dims为 (2, 0, 1)，则输出是 shape 为 [5, 2, 3] 的 tensor。

其本质是一个零拷贝语义的轴重排：数据值本身不变，变化的是逻辑索引到物理存储的映射关系。输出 shape 的第i维等于输入 shape 的第dims[i]维。由此可以推出两条核心规则：

1. dims的个数必须与self的维度数量一致，且每个取值需在[-self 的维度数量, self 的维度数量 - 1]范围内，负数表示从最后一维向前计数；
2. 输出的元素总数与self相同（置换不改变总容量），out的 dtype 必须与self一致，维度最大不超过 8 维。

例如 4D 输入 shape 为[a, b, c, d]，dims = [0, 2, 1, 3]时，输出 shape 为[a, c, b, d]。

3. 参数说明

下表完整继承自 README 的参数说明：

3.1 数据类型支持范围

按产品区分（继承自 README）：

• Atlas A3 训练/推理系列产品：数据类型支持 FLOAT、FLOAT16、BFLOAT16；
• Kirin X90/Kirin 9030 处理器系列产品：数据类型支持 FLOAT、FLOAT16。

源码侧可以印证这一划分。OpDef中x/y的DataType声明为{DT_FLOAT16, DT_BF16, DT_FLOAT}的组合（见 transpose_v2_def.cpp），而 Kirin 配置中收窄为{DT_FLOAT16, DT_FLOAT}。在 device 侧入口 transpose_v2.cpp 中，BF16 分支带有架构判断：

#if ORIG_DTYPE_X == DT_FLOAT16 || (!(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3003 || __NPU_ARCH__ == 3113)) && ORIG_DTYPE_X == DT_BF16)

即当 NPU 架构编号为 3003/3113（Kirin 系列 NPU）时，BF16 编译分支被排除，只有 FLOAT16 走Transpose021<half>/Transpose102<half>路径，FLOAT32 走独立分支。这与“Kirin 平台不支持 BF16”的产品约束在实现层面完全一致。

4. 调用方式：通过 aclnnPermute 接口调用 TransposeV2

README 的调用说明表给出唯一的调用方式：

这里有一个关键的仓库级事实：TransposeV2 并不单独导出自己的 aclnn 接口，而是复用Transpose算子的aclnnPermute两段式接口。从构建配置可以直接确认：

• conversion/transpose_v2/CMakeLists.txt 中声明add_all_modules_sources(OPTYPE transpose_v2 ACLNNTYPE aclnn_exclude DEPENDENCIES transpose)，即本算子的 aclnn 头文件被排除（aclnn_exclude），并显式依赖transpose模块；
• conversion/transpose/CMakeLists.txt 中transpose模块同样标记ACLNNTYPE aclnn_exclude并带有DISABLE_IN_OPP TRUE，说明该接口由 CANN 侧aclnnop包统一提供，接口文档见 aclnnPermute.md。

因此实际调用代码#include "aclnnop/aclnn_permute.h"，调用aclnnPermuteGetWorkspaceSize/aclnnPermute两个函数。aclnnPermute的两段式函数原型为：

aclnnStatus aclnnPermuteGetWorkspaceSize( const aclTensor* self, const aclIntArray* dims, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor); aclnnStatus aclnnPermute( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream);

第一段接口完成入参校验并返回 workspace 大小与执行器；若校验失败会返回ACLNN_ERR_PARAM_NULLPTR（161001，self/dims/out 空指针）或ACLNN_ERR_PARAM_INVALID（161002，dtype 不支持、self 与 out 类型不一致、超过 8 维、dims 越界等）。需要说明的是，通用aclnnPermute接口对 dtype 的声明范围较宽，而 TransposeV2 算子实际支持的数据类型以第 3 节表格为准。

4.1 完整示例代码

仓库提供了可直接参考的完整样例 test_aclnn_transpose_v2.cpp，以selfshape 为{3, 3, 2}、dims = {0, 2, 1}、输出 shape 为{3, 2, 3}的 FLOAT32 张量为例，完整流程如下（核心节选）：

#include <iostream> #include <vector> #include "acl/acl.h" #include "aclnnop/aclnn_permute.h" #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vector<int64_t>& shape) { int64_t shapeSize = 1; for (auto i : shape) { shapeSize *= i; } return shapeSize; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法，初始化 auto ret = aclInit(nullptr); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); ret = aclrtSetDevice(deviceId); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); ret = aclrtCreateStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); return 0; } template <typename T> int CreateAclTensor( const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size = GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); // 计算连续tensor的strides std::vector<int64_t> strides(shape.size(), 1); for (int64_t i = shape.size() - 2; i >= 0; i--) { strides[i] = shape[i + 1] * strides[i + 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor = aclCreateTensor( shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { // 1. （固定写法）device/stream初始化，参考acl API文档 // 根据自己的实际device填写deviceId int32_t deviceId = 0; aclrtStream stream; auto ret = Init(deviceId, &stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); // 2. 构造输入与输出，需要根据API的接口自定义构造 std::vector<int64_t> selfShape = {3, 3, 2}; std::vector<int64_t> dimsData = {0, 2, 1}; std::vector<int64_t> outShape = {3, 2, 3}; void* selfDeviceAddr = nullptr; void* outDeviceAddr = nullptr; aclTensor* self = nullptr; aclTensor* out = nullptr; std::vector<float> selfHostData = {1, 1, 2, 2, 3, 3, 4, 4, 5, 5, 6, 6, 7, 7, 8, 8, 9, 9}; std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; // 创建self aclTensor ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self); CHECK_RET(ret == ACL_SUCCESS, return ret); // 创建dims aclIntArray aclIntArray* dims = aclCreateIntArray(dimsData.data(), dimsData.size()); // 创建out aclTensor ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out); CHECK_RET(ret == ACL_SUCCESS, return ret); // 3. 调用CANN算子库API，需要修改为具体的Api名称 uint64_t workspaceSize = 0; aclOpExecutor* executor; // 调用aclnnPermute第一段接口 ret = aclnnPermuteGetWorkspaceSize(self, dims, out, &workspaceSize, &executor); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnPermuteGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr = nullptr; if (workspaceSize > static_cast<uint64_t>(0)) { ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); } // 调用aclnnPermute第二段接口 ret = aclnnPermute(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnPermute failed. ERROR: %d\n", ret); return ret); // 4. （固定写法）同步等待任务执行结束 ret = aclrtSynchronizeStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧 auto size = GetShapeSize(outShape); std::vector<float> resultData(size, 0); ret = aclrtMemcpy( resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); for (int64_t i = 0; i < size; i++) { LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]); } // 6. 释放aclTensor和aclIntArray aclDestroyTensor(self); aclDestroyIntArray(dims); aclDestroyTensor(out); // 7. 释放device 资源 aclrtFree(selfDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize > static_cast<uint64_t>(0)) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }

示例的关键实操要点：

1. 构造 aclTensor 时必须传 strides：示例中用for循环从右向左计算连续张量的 strides 后交给aclCreateTensor，数据格式固定为ACL_FORMAT_ND；
2. dims 用 aclIntArray 承载：aclCreateIntArray(dimsData.data(), dimsData.size())，与算子定义中perm输入为 INT32/INT64 数组相符；
3. workspace 按需申请：只有workspaceSize > 0时才调用aclrtMalloc申请，否则以空指针调用第二段接口即可；
4. 结果校验：同步后把 device 侧out拷回 host 打印，可以直观验证置换后的数据布局。

5. 源码级实现剖析：Tiling 策略与 Kernel 分发

README 的“约束说明”标注为“无”，但从源码结构看，TransposeV2 的 Tiling 实现针对特定 perm 模式做了高度特化的分核分块策略，理解这部分有助于评估算子在不同 shape 下的行为。

5.1 Tiling 入口：perm 白名单分发

Host 侧 Tiling 入口在 transpose_v2_tiling.cpp 的Tiling4TransposeV2中。它先根据perm输入的数据类型（INT32 或 INT64）解析出轴序数组，然后做白名单分发：

if (perm == std::vector<int64_t>{0, 2, 1}) { ret = DoOpTiling<Transpose021Tiling>(context); } else if (perm == std::vector<int64_t>{1, 0, 2} || perm == std::vector<int64_t>{0, 2, 1, 3}) { ret = DoOpTiling<Transpose102Tiling>(context); } else { OP_LOGE(context->GetNodeName(), "Unsupported perm."); return ge::GRAPH_FAILED; }

从源码结构看，当前 Tiling 实现覆盖三类典型置换：

• {0, 2, 1}：交换后两轴，对应Transpose021Tiling；
• {1, 0, 2}：交换前两轴，对应Transpose102Tiling；
• {0, 2, 1, 3}：4D 交换中间两轴，复用Transpose102Tiling（在 tiling 中按“合轴后都是 0213”的思路处理）。

其他 perm 值会记录Unsupported perm.错误并返回GRAPH_FAILED。最终 Tiling 数据通过IMPL_OP_OPTILING(TransposeV2)宏统一注册，并声明.TilingInputsDataDependency({1})，表明 Tiling 依赖第 2 个输入（perm 数组）的取值（见 transpose_v2_tiling.cpp）。

5.2 Tiling key 设计与 kernel 分发

Tiling 计算完成后会生成一个tilingKey，device 侧 kernel 据此选择具体执行分支。两个 Tiling 类的 key 生成逻辑如下：

021 模式（transpose021_tiling.h）：

void ComputeTilingKey() { params.tilingKey = params.typeSize * TYPE_KEY; }

其中TYPE_KEY = 10，故 FLOAT16/BF16（typeSize=2）得到 key=20，FLOAT32（typeSize=4）得到 key=40。

102/0213 模式（transpose102_tiling.h）：

// permSize 为 3 时基础值为 PERM_KEY(100)，为 4（即 0213）时为基础值*2（200） params.tilingKey += params.typeSize * TYPE_KEY; // dim3Len >= block*4 或 perm 为 4 维时，追加 COPY_KEY(1) 表示走纯 copy 模式 params.tilingKey += COPY_KEY;

组合出的 key 与 transpose_v2.cpp 中的 kernel 分发一一对应：

kernel 入口函数为：

extern "C" __global__ __aicore__ void transpose_v2( GM_ADDR x, GM_ADDR perm, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)

通过GET_TILING_DATA读取 host 计算好的 Tiling 数据，再以TILING_KEY_IS(20)等宏做分支选择，把控制权交给Transpose021、Transpose102、Transpose0213三个执行类。

5.3 分核与分块策略

分核（sub-core）：021 模式把前dimNum-2维折叠为 batch 数inputNC，再按 AIV 核数均分（见 transpose021_tiling.h）：

params.tasksPerCore = params.inputNC / params.coreNum; params.tasksTail = params.inputNC % params.coreNum;

尾部若干核（tasksTail个）每核多领一个任务，保证负载均衡。102 模式则根据dim1Len >= dim2Len选择subMode（0 或 1），决定沿哪个维度切分任务，且当dim0Len != 1时按dim0Len分核（见 transpose102_tiling.h）。

分块（sub-block）：021 模式把后两轴视为 H/W，按TRANS_BLOCK = 16对齐 H、按BLOCK_SIZE / typeSize（BLOCK_SIZE=32，即 FP16 下 16 个元素、FP32 下 8 个元素）对齐 W，并要求对齐后不超过LIMIT_H = 128、LIMIT_W = 128，否则报Unsupported shape.错误（见 transpose021_tiling.h）。随后按 UB 大小（BUFFER_NUM = 4）反推单核一次可搬运的行数hOnceMax，若能容纳则启用双缓冲（doubleBuffer = 2），否则退化为单缓冲。

这些常数统一定义在 transpose_v2_tiling.h：

constexpr uint64_t BLOCK_SIZE = 32; constexpr uint64_t TRANS_BLOCK = 16; constexpr uint64_t NO_TRANS_KEY = 1000000; constexpr uint64_t TYPE_KEY = 10; constexpr uint64_t LIMIT_H = 128; constexpr uint64_t LIMIT_W = 128; constexpr uint64_t BUFFER_NUM = 4; constexpr uint64_t COPY_KEY = 1; constexpr uint64_t PERM_KEY = 100;

device 侧缓冲：kernel 常量定义在 transpose_v2.h，其中MAX_UB_SIZE = 192 * 1024 - 256。以Transpose021为例（transpose021.h），每个核先根据blockIdx、tasksPerCore、tasksTail计算自己负责的全局起始偏移，再用TPipe::InitBuffer(queIn/queOut, doubleBuffer, MAX_UB_SIZE / 2 / doubleBuffer)划分输入/输出双缓冲队列，最后按 H/W 是否对齐选择ProcessOp<ISWALIGN, ISHALIGN>模板实例执行“CopyIn → Compute（Trans）→ CopyOut”流水线。Transpose102与Transpose0213的结构类似，前者通过subMode决定源/目的指针的 stride 计算方式（见 transpose102.h），后者针对 4D 的{0,2,1,3}场景实现 copy 模式。

5.4 单元测试验证

仓库为该算子提供了两层 UT：

• Tiling 测试：test_transpose_v2_tiling.cpp 以 FLOAT16、shape{1, 30, 68}、perm{0, 2, 1}为例，断言 tilingKey 为 20、workspace 为 16777216，并校验完整的 tilingData 序列化串（其中包含inputH16Align=32、inputWAlign=80、hOnce=256等由第 5.3 节算法推出的中间量），是对分块公式的精确回归；
• Kernel 测试：test_transpose_v2.cpp 在 device 侧对比输出张量数据，验证置换结果的正确性。

6. 约束说明与使用注意事项

README 的“约束说明”章节原文为：无。在此之上，结合源码可补充以下工程实践层面的注意事项：

1. perm 模式覆盖范围：从源码结构看，当前 Tiling 实现按 perm 白名单（{0,2,1}、{1,0,2}、{0,2,1,3}）特化，其余排列在 Tiling 阶段返回GRAPH_FAILED并记录Unsupported perm.。在 A2/A3/Kirin 平台使用前建议确认目标 perm 属于受支持的模式；
2. shape 限制：021 路径要求对齐后的 H/W 不超过 128，超限会报Unsupported shape.；输入维度需大于等于 2；
3. 数据类型：A3 平台支持 FLOAT/FLOAT16/BFLOAT16，Kirin 平台仅支持 FLOAT/FLOAT16，且out的 dtype 必须与self一致；dims支持 INT32/INT64 两种整型数组；
4. 维度上限：self/out最大 8 维；
5. 确定性计算：根据 aclnnPermute.md 的约束说明，aclnnPermute默认确定性实现，即 TransposeV2 经该接口调用时输出是确定的；
6. 接口归属：调用侧应包含aclnnop/aclnn_permute.h并使用aclnnPermute两段式接口，而不是寻找aclnnTransposeV2这类独立入口——本算子通过ACLNNTYPE aclnn_exclude DEPENDENCIES transpose复用 Transpose 模块的 aclnn 接口。

7. 核心文件索引

通过以上链路可以完整回答三个问题：TransposeV2 在哪些产品上可用（第 1 节）、如何正确调用它（第 4 节）、以及它在 NPU 内部如何通过 Tiling key 把“分核—分块—双缓冲流水线”串起来执行维度置换（第 5 节）。

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