Star869
1092
代码
代码
Issues71
Pull Requests223
Actions
项目讨论
Wiki
分析
Star869
1092
已合并
support tensor api reg vector compute select/fill/deinterleave api #5510
li-xingyue-lxy创建于 19 天前
support tensor api reg vector compute select/fill/deinterleave api #5510
已合并
li-xingyue-lxy创建于 19 天前
从li-xingyue-lxy/asc-devkitmaster合入到cann/asc-devkitmaster
共 32 个文件变更+1859-2
Mdocs/zh/api/README.md+7-0
@@ -1443,6 +1443,13 @@
1443 - [L1 Buffer到BiasTable Buffer数据搬运](SIMD-API/tensor_api/data_move/l1_bias_table_data_move.md)1443 - [L1 Buffer到BiasTable Buffer数据搬运](SIMD-API/tensor_api/data_move/l1_bias_table_data_move.md)
1444 - [Unified Buffer内数据搬运](SIMD-API/tensor_api/data_move/ub_ub_data_move/ub_ub_data_move.md)1444 - [Unified Buffer内数据搬运](SIMD-API/tensor_api/data_move/ub_ub_data_move/ub_ub_data_move.md)
1445 - [copy（Unified Buffer内数据搬运）](SIMD-API/tensor_api/data_move/ub_ub_data_move/copy_ub_to_ub.md)1445 - [copy（Unified Buffer内数据搬运）](SIMD-API/tensor_api/data_move/ub_ub_data_move/copy_ub_to_ub.md)
1446+ - [reg矢量计算](SIMD-API/tensor_api/reg_vector_compute/reg_vector_compute.md)
1447+ - [选择与比较](SIMD-API/tensor_api/reg_vector_compute/compare_and_select/compare_and_select.md)
1448+ - [select](SIMD-API/tensor_api/reg_vector_compute/compare_and_select/select.md)
1449+ - [数据填充](SIMD-API/tensor_api/reg_vector_compute/data_padding/data_padding.md)
1450+ - [fill](SIMD-API/tensor_api/reg_vector_compute/data_padding/fill.md)
1451+ - [数据重排](SIMD-API/tensor_api/reg_vector_compute/data_reorder/data_reorder.md)
1452+ - [deinterleave](SIMD-API/tensor_api/reg_vector_compute/data_reorder/deinterleave.md)
1446 - [工具接口](SIMD-API/tensor_api/utils/tool_functions.md)1453 - [工具接口](SIMD-API/tensor_api/utils/tool_functions.md)
1447 - [编译期整型别名](SIMD-API/tensor_api/utils/compile_time_integer_aliases.md)1454 - [编译期整型别名](SIMD-API/tensor_api/utils/compile_time_integer_aliases.md)
1448 - [维度保留标记](SIMD-API/tensor_api/utils/dimension_keep_mark.md)1455 - [维度保留标记](SIMD-API/tensor_api/utils/dimension_keep_mark.md)
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/compare_and_select/compare_and_select.md+5-0
@@ -0,0 +1,5 @@
1+# 概述
2+ 
3+# 比较与选择
4+ 
5+- **[select](select.md)**
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/compare_and_select/select.md+80-0
@@ -0,0 +1,80 @@
1+# select
2+ 
3+## 产品支持情况
M
Mmunanhw18 天前

新增文件需要修改api目录下的README.md索引文件

likedislike
4+ 
5+<!-- npu="950" id1 -->
6+- Ascend 950PR/Ascend 950DT：支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- Atlas 200I/500 A2 推理产品：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- Atlas 推理系列产品 AI Core：不支持
19+<!-- end id5 -->
20+<!-- npu="310p" id6 -->
21+- Atlas 推理系列产品 Vector Core：不支持
22+<!-- end id6 -->
23+<!-- npu="910" id7 -->
24+- Atlas 训练系列产品：不支持
25+<!-- end id7 -->
26+ 
27+## 功能说明
28+ 
29+头文件路径：`"include/tensor_api/experimental/arch/vector/compare_and_select.h"`。
30+ 
31+给定两个源操作数src0和src1，根据condition的比特位值选取元素，得到目的操作数。当condition的比特位是1时，从src0中选取对应位置的数，比特位是0时从src1选取对应位置的数。
32+ 
33+## 函数原型
34+ 
35+```cpp
36+template <typename T>
37+__simd_callee__ inline reg_tensor<T> select(reg_tensor<bool> condition, reg_tensor<T> src0, reg_tensor<T> src1)
38+```
39+ 
40+## 参数说明
41+ 
42+| 参数 | 输入/输出 | 描述 |
43+| --- | --- | --- |
44+| condition | 输入 | 选择条件，类型为 `reg_tensor<bool>`。每个元素决定对应位置选择 `src0` 还是 `src1`。 |
45+| src0 | 输入 | 条件为 `true` 时选择的源操作数，类型为 `reg_tensor<T>`。其 `mask` 作为返回值的有效元素掩码。 |
46+| src1 | 输入 | 条件为 `false` 时选择的源操作数，类型为 `reg_tensor<T>`。 |
47+ 
48+## 返回值说明
49+ 
50+返回 `reg_tensor<T>`，包含逐元素选择的结果。返回值的 `mask` 与 `src0.mask` 相同。
51+ 
52+## 数据类型
53+ 
54+支持的数据类型为：bool、int8_t、uint8_t、fp8_e4m3fn_t、hifloat8_t、fp8_e5m2_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float。
55+ 
56+## 约束说明
57+ 
58+通过src0的mask参数控制的未选中的元素是未定义行为。
59+ 
60+## 调用示例
61+ 
62+```cpp
63+template <typename T, typename U, typename V>
64+__simd_vf__ inline void select_vf(T x_tensor, U y_tensor, V z_tensor,
65+ uint32_t repeat_times, uint32_t one_repeat_size)
66+{
67+ uint32_t count = TOTAL_LENGTH;
68+ using data_type = typename T::data_type;
B
Bbianque18 天前

data_type是否要和上面的typename 风格保持一致？

likedislike
69+ auto condition = asc::te::experimental::make_mask<asc::te::experimental::mask_pattern::vl8, data_type>();
70+ for (uint16_t i = 0; i < repeat_times; ++i) {
71+ auto mask = asc::te::experimental::update_mask<data_type>(count);
72+ uint32_t offset = i * one_repeat_size;
73+ const auto coord = asc::te::make_coord(offset);
74+ auto x_reg = x_tensor.load(coord).with_mask(mask);
75+ auto y_reg = y_tensor.load(coord).with_mask(mask);
76+ auto z_reg = asc::te::experimental::select((x_reg > y_reg), x_reg, y_reg);
77+ z_tensor.store(coord, z_reg);
78+ }
79+}
80+```
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/data_padding/data_padding.md+3-0
@@ -0,0 +1,3 @@
1+# 数据填充
2+ 
3+- **[fill](fill.md)**
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/data_padding/fill.md+118-0
@@ -0,0 +1,118 @@
1+# fill
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- Ascend 950PR/Ascend 950DT：支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- Atlas 200I/500 A2 推理产品：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- Atlas 推理系列产品AI Core：不支持
19+<!-- end id5 -->
20+<!-- npu="310p" id6 -->
21+- Atlas 推理系列产品Vector Core：不支持
22+<!-- end id6 -->
23+<!-- npu="910" id7 -->
24+- Atlas 训练系列产品：不支持
25+<!-- end id7 -->
26+ 
27+## 功能说明
28+ 
29+头文件路径：`"include/tensor_api/experimental/arch/vector/compare_and_select.h"`。
30+ 
31+支持scalar和tensor两种模式：
32+ 
33+- scalar模式：将value广播到寄存器，并保存在目的寄存器中（如果有mask，则保存在目的寄存器中被mask筛选的位置）。
34+- tensor模式：将src的最低位元素广播到寄存器，并保存在目的寄存器中被mask筛选的位置。
35+ 
36+## 函数原型
37+ 
38+- 将value广播到目的寄存器
39+ 
40+ ```cpp
41+ template <size_t N = 1, typename T>
42+ __simd_callee__ inline reg_tensor<T> fill(const T& value, const reg_tensor<bool>& mask = all_mask<T>());
43+ ```
44+ 
45+- 将寄存器src的最低位元素广播到目的寄存器
46+ 
47+ ```cpp
48+ template <size_t N = 1, typename T>
49+ __simd_callee__ inline reg_tensor<T> fill(const reg_tensor<T>& src, const reg_tensor<bool>& mask = all_mask<T>());
50+ ```
51+ 
52+## 参数说明
53+ 
54+**表1** 模板参数说明
55+ 
56+| 参数名 | 描述 |
57+| --- | --- |
58+| T | 操作数数据类型。支持的数据类型请参考[数据类型](#数据类型)。 |
59+ 
60+**表2** 函数参数说明
61+ 
62+| 参数名 | 输入/输出 | 描述 |
63+| --- | --- | --- |
64+| value | 输入 | 源操作数。<br>类型为标量。 |
65+| src | 输入 | 源操作数。<br>类型为[reg_tensor]。 |
66+| mask | 输入 | 源操作数中元素有效性的指示。 |
67+ 
68+## 数据类型
69+ 
70+支持的数据类型为：int8_t、uint8_t、fp8_e8m0_t、fp8_e5m2_t、fp8_e4m3fn_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float。
71+ 
72+ 
73+## 返回值说明
74+ 
75+返回 `reg_tensor<T>`，包含数据广播后的结果。返回值的 `mask` 与 输入的 `mask` 相同。
76+ 
77+## 约束说明
78+ 
79+mask仅筛选dst中写入广播值的位置，不筛选src中的元素。无论mask的最低位是否有效，待广播的数据均为src的最低位元素。
80+ 
81+## 调用示例
82+ 
83+- 示例一
84+ 
85+ ```cpp
86+ template <typename T, typename U>
87+ __simd_vf__ inline void fill_vf(T y_tensor, U scalar, uint32_t repeat_times, uint32_t one_repeat_size)
88+ {
89+ uint32_t count = total_length;
90+ using data_type = typename T::data_type;
91+ for (uint16_t i = 0; i < repeat_times; ++i) {
92+ auto mask = asc::te::experimental::update_mask<data_type>(count);
93+ uint32_t offset = i * one_repeat_size;
94+ const auto coord = asc::te::make_coord(offset);
95+ auto y_reg = asc::te::experimental::fill(scalar, mask);
96+ asc::te::experimental::store(y_tensor, coord, y_reg);
97+ }
98+ }
99+ ```
100+ 
101+- 示例二
102+ 
103+ ```cpp
104+ template <typename T, typename U>
105+ __simd_vf__ inline void fill_vf(T x_tensor, U y_tensor, uint32_t repeat_times, uint32_t one_repeat_size)
106+ {
107+ uint32_t count = total_length;
108+ using data_type = typename T::data_type;
109+ for (uint16_t i = 0; i < repeat_times; ++i) {
110+ auto mask = asc::te::experimental::update_mask<data_type>(count);
111+ uint32_t offset = i * one_repeat_size;
112+ const auto coord = asc::te::make_coord(offset);
113+ auto x_reg = asc::te::experimental::load(x_tensor, coord);
114+ auto y_reg = asc::te::experimental::fill(x_reg, mask);
115+ asc::te::experimental::store(y_tensor, coord, y_reg);
116+ }
117+ }
118+ ```
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/data_reorder/data_reorder.md+3-0
@@ -0,0 +1,3 @@
1+# 数据重排
2+ 
3+- **[deinterleave](deinterleave.md)**
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/data_reorder/deinterleave.md+87-0
@@ -0,0 +1,87 @@
1+# deinterleave
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- Ascend 950PR/Ascend 950DT：支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- Atlas 200I/500 A2 推理产品：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- Atlas 推理系列产品AI Core：不支持
19+<!-- end id5 -->
20+<!-- npu="310p" id6 -->
21+- Atlas 推理系列产品Vector Core：不支持
22+<!-- end id6 -->
23+<!-- npu="910" id7 -->
24+- Atlas 训练系列产品：不支持
25+<!-- end id7 -->
26+ 
27+## 功能说明
28+ 
29+头文件路径：`"include/tensor_api/experimental/arch/vector/data_reorder.h"`。
30+ 
31+给定源操作数寄存器src0和src1，将src0和src1中的元素解交织存入结果操作数dst0和dst1中。解交织排列方式如下图所示，其中每个方格代表一个元素：
32+ 
33+![DeInterleave示意图](../../../../figures/reg_regdeinterleave.png)
34+ 
35+## 函数原型
36+ 
37+```cpp
38+template <typename T>
39+__simd_callee__ inline reg_pair<T> deinterleave(reg_tensor<T> src0, reg_tensor<T> src1);
40+```
41+ 
42+## 参数说明
43+ 
44+**表1** 模板参数说明
45+ 
46+| 参数名 | 描述 |
47+| --- | --- |
48+| T | 目的操作数和源操作数的数据类型。支持的数据类型请参考[数据类型](#数据类型)。 |
49+ 
50+**表2** 函数参数说明
51+ 
52+| 参数名 | 输入/输出 | 描述 |
53+| --- | --- | --- |
54+| src0<br>src1 | 输入 | 源操作数。<br>类型为[reg_tensor]。 |
M
Mmunanhw18 天前

reg_tensor后续会补充链接么

likedislike
L
li-xingyue-lxy
18 天前 评论：
55+ 
56+## 数据类型
57+ 
58+支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float。
B
Bbianque18 天前

下划线，是否要加转义？需要统一

likedislike
59+ 
60+## 返回值说明
61+ 
62+返回 `reg_pair<T>`，由两个 `reg_tensor<T>` 组成，存放交织后的两个结果值。返回值的 `mask` 是和数据类型T相关的满mask。
63+ 
64+## 约束说明
65+ 
66+- src0、src1的数据类型需要保持一致。
67+- src0和src1可以为同一个reg_tensor。
68+ 
69+## 调用示例
70+ 
71+```cpp
72+template <typename Src0TensorType, typename Src1TensorType, typename Dst0TensorType, typename Dst1TensorType>
73+__simd_vf__ inline void data_reorder_vf(
74+ Src0TensorType src0_tensor, Src1TensorType src1_tensor, Dst0TensorType dst0_tensor, Dst1TensorType dst1_tensor,
75+ uint32_t repeat_times, uint32_t one_repeat_size)
76+{
77+ for (uint16_t i = 0; i < repeat_times; ++i) {
78+ uint32_t offset = i * one_repeat_size;
79+ const auto coord = asc::te::make_coord(offset);
80+ auto src0_reg = asc::te::experimental::load(src0_tensor, coord);
81+ auto src1_reg = asc::te::experimental::load(src1_tensor, coord);
82+ auto result = asc::te::experimental::deinterleave(src0_reg, src1_reg);
83+ asc::te::experimental::store(dst0_tensor, coord, result.first);
84+ asc::te::experimental::store(dst1_tensor, coord, result.second);
85+ }
86+}
87+```
Adocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/reg_vector_compute.md+5-0
@@ -0,0 +1,5 @@
1+# Reg矢量计算
2+ 
3+- **[选择与比较](compare_and_select/compare_and_select.md)**
4+- **[数据填充](data_padding/data_padding.md)**
5+- **[数据重排](data_reorder/data_reorder.md)**
Mdocs/zh/api/SIMD-API/tensor_api/tensor_api.md+1-0
@@ -28,6 +28,7 @@ Tensor API提供基于Tensor的数据搬运、矩阵计算和矢量计算能力
28- [Algorithm](algorithm.md)：介绍数据搬运和矩阵计算中Operation、Trait与Atom的作用、组合关系及使用方式。28- [Algorithm](algorithm.md)：介绍数据搬运和矩阵计算中Operation、Trait与Atom的作用、组合关系及使用方式。
29- [矩阵计算](matrix_compute/matrix_compute.md)：介绍矩阵数据搬入、矩阵乘加、矩阵结果搬出和相关关键特性。29- [矩阵计算](matrix_compute/matrix_compute.md)：介绍矩阵数据搬入、矩阵乘加、矩阵结果搬出和相关关键特性。
30- [数据搬运](data_move/data_move.md)：介绍多种数据搬运通路，包括Global Memory与Unified Buffer之间的双向搬运，Global Memory到L1 Buffer的搬运等。30- [数据搬运](data_move/data_move.md)：介绍多种数据搬运通路，包括Global Memory与Unified Buffer之间的双向搬运，Global Memory到L1 Buffer的搬运等。
31+- [Reg矢量计算](reg_vector_compute/reg_vector_compute.md)：介绍reg矢量计算相关的接口。
31- [工具接口](utils/tool_functions.md)：介绍编译期整数、维度保留标记和Layout Pattern查询接口，以及Layout、Tensor、Coord、Shape和量化参数的类型判断等工具接口的使用方式。32- [工具接口](utils/tool_functions.md)：介绍编译期整数、维度保留标记和Layout Pattern查询接口，以及Layout、Tensor、Coord、Shape和量化参数的类型判断等工具接口的使用方式。
32- [废弃接口](deprecated_interface/deprecated_interface.md)：介绍用于兼容的大驼峰风格接口。33- [废弃接口](deprecated_interface/deprecated_interface.md)：介绍用于兼容的大驼峰风格接口。
33 34 
Mdocs/zh/api/SIMD-API/tensor_api/tensor_api_list.md+12-2
@@ -94,9 +94,19 @@
94| [Batch搬运](matrix_compute/cube_store_key_features/batch_copy.md) | 介绍矩阵结果的多Batch搬出能力。 |94| [Batch搬运](matrix_compute/cube_store_key_features/batch_copy.md) | 介绍矩阵结果的多Batch搬出能力。 |
95| [L0C到Unified Buffer双目标模式](matrix_compute/cube_store_key_features/l0c_to_ub_dual_dst.md) | 介绍L0C到Unified Buffer的双目标搬出模式。 |95| [L0C到Unified Buffer双目标模式](matrix_compute/cube_store_key_features/l0c_to_ub_dual_dst.md) | 介绍L0C到Unified Buffer的双目标搬出模式。 |
96 96 
97+## 矢量计算接口
98+ 
99+**表6** 矢量计算接口列表
100+ 
101+| 接口或类型名 | 功能描述 |
102+| --- | --- |
103+| [选择与比较](reg_vector_compute/compare_and_select/compare_and_select.md) | 根据mask的比特位值，从源操作数src0、src1中选择元素，得到目的操作数。 |
104+| [数据填充](reg_vector_compute/data_padding/data_padding.md) | 根据mask将源操作数src的最低位元素或者一个scalar操作数填充到目的操作数。 |
105+| [数据重排](reg_vector_compute/data_reorder/data_reorder.md) | 给定源操作数src0和src1，将src0和src1中的元素解交织存入结果操作数。 |
106+ 
97## 工具接口107## 工具接口
98 108 
99-**表6** 工具接口列表109+**表7** 工具接口列表
100 110 
101| 接口或类型名 | 功能描述 |111| 接口或类型名 | 功能描述 |
102| --- | --- |112| --- | --- |
@@ -107,7 +117,7 @@
107 117 
108## 废弃接口118## 废弃接口
109 119 
110-**表7** 废弃接口列表120+**表8** 废弃接口列表
111 121 
112| 接口名 | 功能描述 |122| 接口名 | 功能描述 |
113| --- | --- |123| --- | --- |
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/data_reorder/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# ----------------------------------------------------------------------------------------------------------
6+cmake_minimum_required(VERSION 3.16)
7+ 
8+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, sim")
9+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
10+option(CANN_ASC_USE_EXPERIMENTAL "Enable experimental CANN ASC APIs" OFF)
11+ 
12+find_package(ASC REQUIRED)
13+ 
14+project(kernel_samples LANGUAGES ASC CXX)
15+ 
16+add_executable(demo
17+ data_reorder.asc
18+)
19+ 
20+target_link_libraries(demo PRIVATE
21+)
22+ 
23+target_compile_options(demo PRIVATE
24+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
25+ $<$<AND:$<COMPILE_LANGUAGE:ASC>,$<BOOL:${CANN_ASC_USE_EXPERIMENTAL}>>:-DCANN_ASC_USE_EXPERIMENTAL>
26+)
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/data_reorder/README.md+86-0
@@ -0,0 +1,86 @@
1+# data_reorder 样例
2+ 
3+## 概述
4+ 
5+本样例基于静态Tensor API编程范式实现数据重排功能，主要调用实验性接口 `asc::te::experimental::deinterleave`。
6+ 
7+## 本样例支持的产品及CANN软件版本
8+ 
9+| 产品 | CANN软件版本 |
10+|------|-------------|
11+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
12+ 
13+## 目录结构
14+ 
15+```text
16+data_reorder
17+├── scripts/gen_data.py
18+├── CMakeLists.txt
19+├── data_reorder.asc
20+├── data_utils.h
21+└── README.md
22+```
23+ 
24+## 样例描述
25+- 样例功能：
26+ - 两个float向量（各256个元素）交织，输出两个float向量
27+- 样例规格：
28+ <table>
29+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
30+ <tr><td rowspan="3" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
31+ <tr><td align="center">src0</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
32+ <tr><td align="center">src1</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
33+ <tr><td rowspan="3" align="center">样例输出</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
34+ <tr><td align="center">dst0</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
35+ <tr><td align="center">dst1</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
36+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="3" align="center">data_reorder</td></tr>
37+ </table>
38+- 样例实现：
39+ data_reorder_vf函数内调用deinterleave接口进行数据重排：
40+ - 使用load将输入从UB读取到src0_reg和src1_reg寄存器中
41+ - 使用deinterleave接口将标量值填充到目的reg_tensor的每个元素
42+ - 使用store将结果写回UB
43+ - 调用实现
44+ 使用内核调用符<<<>>>调用核函数。
45+ 
46+ 
47+## 编译运行
48+在本样例根目录下执行如下步骤，编译并执行样例。
49+- 配置环境变量
50+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/zh/quick_start.md#prepare&install)，配置环境变量，**当前仅支持使用[CANN master](../../../../../../docs/zh/quick_start.md#cann-install)**。
51+ 
52+ > **说明：** `${install_path}` 为CANN包安装目录，未指定安装目录时默认安装至 `/usr/local/Ascend` 下。
53+ 
54+- 样例执行
55+ 
56+ 在本样例目录下执行如下命令。
57+ ```bash
58+ mkdir -p build && cd build; # 创建并进入build目录
59+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # 编译工程（默认npu模式）
60+ python3 ../scripts/gen_data.py # 生成测试真值数据
61+ ./demo # 执行编译生成的可执行程序，执行样例
62+ ```
63+ 
64+ 使用 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
65+ 
66+ 示例如下：
67+ ```bash
68+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # NPU仿真模式
69+ ```
70+ 
71+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
72+ 
73+- 编译选项说明
74+ 
75+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
76+| ---------------------------| -----------------------------| ---------------------------------------------------|
77+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`sim` | 运行模式：NPU 运行、NPU仿真　　　　　　　|
78+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
79+| `CANN_ASC_USE_EXPERIMENTAL` | `ON`（本样例必选）、`OFF`（默认） | 开启实验性ASC接口 |
80+ 
81+- 执行结果
82+ 
83+ 执行结果如下，说明精度对比成功。
84+ ```bash
85+ test pass!
86+ ```
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/data_reorder/README_en.md+90-0
@@ -0,0 +1,90 @@
1+# data_reorder Example
2+ 
3+## Overview
4+ 
5+This example implements data reordering based on the static Tensor API programming model. It primarily calls the experimental `asc::te::experimental::deinterleave` interface.
6+ 
7+## Supported Products and CANN Versions
8+ 
9+| Product | CANN Version |
10+| --- | --- |
11+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
12+ 
13+## Directory Structure
14+ 
15+```text
16+data_reorder
17+├── scripts/gen_data.py
18+├── CMakeLists.txt
19+├── data_reorder.asc
20+├── data_utils.h
21+└── README_en.md
22+```
23+ 
24+## Example Description
25+ 
26+- Function:
27+ - Deinterleaves two float vectors, each containing 256 elements, and outputs two float vectors.
28+- Specifications:
29+ <table>
30+ <tr><td rowspan="1" align="center">Example Type (OpType)</td><td colspan="3" align="center">AIV Example</td></tr>
31+ <tr><td rowspan="3" align="center">Inputs</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
32+ <tr><td align="center">src0</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
33+ <tr><td align="center">src1</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
34+ <tr><td rowspan="3" align="center">Outputs</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
35+ <tr><td align="center">dst0</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
36+ <tr><td align="center">dst1</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
37+ <tr><td rowspan="1" align="center">Kernel Function</td><td colspan="3" align="center">data_reorder</td></tr>
38+ </table>
39+- Implementation:
40+ The `data_reorder_vf` function calls the `deinterleave` interface to reorder data:
41+ - Use `load` to read the inputs from UB into the `src0_reg` and `src1_reg` registers.
42+ - Use `deinterleave` to deinterleave the source data into the destination `reg_tensor` objects.
43+ - Use `store` to write the results back to UB.
44+ - Invoke the kernel function using the `<<<>>>` kernel launch syntax.
45+ 
46+## Build and Run
47+ 
48+Perform the following steps in the example root directory to build and run the example.
49+ 
50+- Configure environment variables
51+ Configure the environment variables according to the CANN development kit [installation instructions](../../../../../../docs/en/quick_start.md#prepare&install). **Currently, only [CANN master](../../../../../../docs/en/quick_start.md#cann-install) is supported.**
52+ 
53+ > **Note:** `${install_path}` is the CANN package installation directory. If no installation directory is specified, the default directory is `/usr/local/Ascend`.
54+ 
55+- Run the example
56+ 
57+ Run the following commands in the example directory:
58+ 
59+ ```bash
60+ mkdir -p build && cd build; # Create and enter the build directory
61+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # Build the project (NPU mode by default)
62+ python3 ../scripts/gen_data.py # Generate test ground truth data
63+ ./demo # Run the generated executable
64+ ```
65+ 
66+ To use NPU simulation mode, add the `-DCMAKE_ASC_RUN_MODE=sim` option.
67+ 
68+ Example:
69+ 
70+ ```bash
71+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # NPU simulation mode
72+ ```
73+ 
74+ > **Note:** Clear the CMake cache before switching build modes. Run `rm CMakeCache.txt` in the `build` directory, and then run CMake again.
75+ 
76+- Build options
77+ 
78+| Option | Values | Description |
79+| --- | --- | --- |
80+| `CMAKE_ASC_RUN_MODE` | `npu` (default), `sim` | Run mode: NPU execution or NPU simulation |
81+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU architecture: `dav-3510` corresponds to Ascend 950PR/Ascend 950DT |
82+| `CANN_ASC_USE_EXPERIMENTAL` | `ON` (required for this example), `OFF` (default) | Enables experimental ASC interfaces |
83+ 
84+- Execution result
85+ 
86+ The following output indicates that the accuracy comparison is successful:
87+ 
88+ ```bash
89+ test pass!
90+ ```
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/data_reorder/data_reorder.asc+176-0
@@ -0,0 +1,176 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <cmath>
12+#include <cstdint>
13+#include <cstdio>
14+#include <vector>
15+ 
16+#include "acl/acl.h"
17+#include "data_utils.h"
18+#include "tensor_api/experimental/vector_compute.h"
19+#include "tensor_api/tensor.h"
20+ 
21+constexpr uint32_t total_length = 256;
22+constexpr uint8_t buffer_lock_id = 0;
23+ 
24+template <typename Src0TensorType, typename Src1TensorType, typename Dst0TensorType, typename Dst1TensorType>
25+__simd_vf__ inline void data_reorder_vf(
26+ Src0TensorType src0_tensor, Src1TensorType src1_tensor, Dst0TensorType dst0_tensor, Dst1TensorType dst1_tensor,
27+ uint32_t repeat_times, uint32_t one_repeat_size)
28+{
29+ for (uint16_t i = 0; i < repeat_times; ++i) {
30+ uint32_t offset = i * one_repeat_size;
31+ const auto coord = asc::te::make_coord(offset);
32+ auto src0_reg = asc::te::experimental::load(src0_tensor, coord);
33+ auto src1_reg = asc::te::experimental::load(src1_tensor, coord);
34+ auto result = asc::te::experimental::deinterleave(src0_reg, src1_reg);
35+ asc::te::experimental::store(dst0_tensor, coord, result.first);
36+ asc::te::experimental::store(dst1_tensor, coord, result.second);
37+ }
38+}
39+ 
40+template <typename DataType>
41+class Kernel {
42+public:
43+ __aicore__ inline void process(
44+ __gm__ DataType* src0, __gm__ DataType* src1, __gm__ DataType* dst0, __gm__ DataType* dst1)
45+ {
46+ auto layout = asc::te::make_layout(asc::te::make_shape(total_length), asc::te::make_stride(1));
47+ auto src0_gm = asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::gm>(src0), layout);
48+ auto src1_gm = asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::gm>(src1), layout);
49+ auto dst0_gm = asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::gm>(dst0), layout);
50+ auto dst1_gm = asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::gm>(dst1), layout);
51+ 
52+ constexpr uint32_t src0_offset = 0;
53+ constexpr uint32_t src1_offset = total_length * sizeof(DataType);
54+ constexpr uint32_t dst0_offset = src1_offset + total_length * sizeof(DataType);
55+ constexpr uint32_t dst1_offset = dst0_offset + total_length * sizeof(DataType);
56+ auto src0_ub =
57+ asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::ub, DataType>(src0_offset), layout);
58+ auto src1_ub =
59+ asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::ub, DataType>(src1_offset), layout);
60+ auto dst0_ub =
61+ asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::ub, DataType>(dst0_offset), layout);
62+ auto dst1_ub =
63+ asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::ub, DataType>(dst1_offset), layout);
64+ 
65+ asc_lock(pipe_t::PIPE_MTE2, buffer_lock_id);
66+ auto copy_in = asc::te::make_copy(asc::te::copy_gm_to_ub{});
67+ asc::te::copy(copy_in, src0_ub, src0_gm);
68+ asc::te::copy(copy_in, src1_ub, src1_gm);
69+ asc_unlock(pipe_t::PIPE_MTE2, buffer_lock_id);
70+ asc_lock(pipe_t::PIPE_V, buffer_lock_id);
71+ constexpr uint32_t one_repeat_size = asc_get_vf_len() / sizeof(DataType);
72+ uint16_t repeat_times =
73+ static_cast<uint16_t>(AscendC::Std::ceil_align(total_length, one_repeat_size) / one_repeat_size);
74+ data_reorder_vf(src0_ub, src1_ub, dst0_ub, dst1_ub, repeat_times, one_repeat_size);
75+ asc_unlock(pipe_t::PIPE_V, buffer_lock_id);
76+ asc_lock(pipe_t::PIPE_MTE3, buffer_lock_id);
77+ auto copy_out = asc::te::make_copy(asc::te::copy_ub_to_gm{});
78+ asc::te::copy(copy_out, dst0_gm, dst0_ub);
79+ asc::te::copy(copy_out, dst1_gm, dst1_ub);
80+ asc_unlock(pipe_t::PIPE_MTE3, buffer_lock_id);
81+ }
82+};
83+ 
84+template <typename DataType>
85+__global__ __vector__ void data_reorder(
86+ __gm__ uint8_t* src0, __gm__ uint8_t* src1, __gm__ uint8_t* dst0, __gm__ uint8_t* dst1)
87+{
88+ Kernel<DataType> kernel;
89+ kernel.process(
90+ reinterpret_cast<__gm__ DataType*>(src0), reinterpret_cast<__gm__ DataType*>(src1),
91+ reinterpret_cast<__gm__ DataType*>(dst0), reinterpret_cast<__gm__ DataType*>(dst1));
92+}
93+ 
94+static bool CompareResult(const void* outputData, const void* goldenData, uint32_t outSize)
95+{
96+ constexpr float eps = 1e-4;
97+ int64_t wrongNum = 0;
98+ 
99+ for (size_t i = 0; i < outSize / sizeof(float); i++) {
100+ float a = (reinterpret_cast<const float*>(outputData))[i];
101+ float b = (reinterpret_cast<const float*>(goldenData))[i];
102+ float ae = std::abs(a - b);
103+ float re = ae / std::abs(b);
104+ if (ae > eps && re > eps) {
105+ printf("CompareResult golden.bin failed at index %zu: output is %lf, golden is %lf\n", i, a, b);
106+ wrongNum++;
107+ }
108+ }
109+ if (wrongNum != 0) {
110+ return false;
111+ } else {
112+ printf("CompareResult success!\n");
113+ return true;
114+ }
115+}
116+ 
117+int32_t main(int32_t argc, char* argv[])
118+{
119+ aclInit(nullptr);
120+ int32_t deviceId = 0;
121+ aclrtSetDevice(deviceId);
122+ aclrtStream stream = nullptr;
123+ aclrtCreateStream(&stream);
124+ uint32_t numBlocks = 1;
125+ size_t byteSize = total_length * sizeof(float);
126+ uint8_t *src0Host, *src1Host, *dst0Host, *dst1Host;
127+ uint8_t *src0Device, *src1Device, *dst0Device, *dst1Device;
128+ aclrtMallocHost((void**)(&src0Host), byteSize);
129+ aclrtMallocHost((void**)(&src1Host), byteSize);
130+ aclrtMallocHost((void**)(&dst0Host), byteSize);
131+ aclrtMallocHost((void**)(&dst1Host), byteSize);
132+ aclrtMalloc((void**)&src0Device, byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
133+ aclrtMalloc((void**)&src1Device, byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
134+ aclrtMalloc((void**)&dst0Device, byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
135+ aclrtMalloc((void**)&dst1Device, byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
136+ size_t fileSize = byteSize;
137+ ReadFile("./input/input_src0.bin", fileSize, src0Host, byteSize);
138+ ReadFile("./input/input_src1.bin", fileSize, src1Host, byteSize);
139+ aclrtMemcpy(src0Device, byteSize, src0Host, byteSize, ACL_MEMCPY_HOST_TO_DEVICE);
140+ aclrtMemcpy(src1Device, byteSize, src1Host, byteSize, ACL_MEMCPY_HOST_TO_DEVICE);
141+ data_reorder<float><<<numBlocks, 0, stream>>>(src0Device, src1Device, dst0Device, dst1Device);
142+ aclrtSynchronizeStream(stream);
143+ aclrtMemcpy(dst0Host, byteSize, dst0Device, byteSize, ACL_MEMCPY_DEVICE_TO_HOST);
144+ aclrtMemcpy(dst1Host, byteSize, dst1Device, byteSize, ACL_MEMCPY_DEVICE_TO_HOST);
145+ WriteFile("./output/output_dst0.bin", dst0Host, byteSize);
146+ WriteFile("./output/output_dst1.bin", dst1Host, byteSize);
147+ void* g0;
148+ aclrtMallocHost((void**)(&g0), byteSize);
149+ size_t gs0 = byteSize;
150+ ReadFile("./output/golden_dst0.bin", gs0, g0, byteSize);
151+ bool r0 = CompareResult(dst0Host, g0, byteSize);
152+ aclrtFreeHost(g0);
153+ void* g1;
B
Bbianque18 天前

命名风格排查可理解性

likedislike
154+ aclrtMallocHost((void**)(&g1), byteSize);
155+ size_t gs1 = byteSize;
156+ ReadFile("./output/golden_dst1.bin", gs1, g1, byteSize);
157+ bool r1 = CompareResult(dst1Host, g1, byteSize);
158+ aclrtFreeHost(g1);
159+ if (r0 && r1) {
160+ printf("test pass!\n");
161+ } else {
162+ printf("test failed!\n");
163+ }
164+ aclrtFree(src0Device);
165+ aclrtFree(src1Device);
166+ aclrtFree(dst0Device);
167+ aclrtFree(dst1Device);
168+ aclrtFreeHost(src0Host);
169+ aclrtFreeHost(src1Host);
170+ aclrtFreeHost(dst0Host);
171+ aclrtFreeHost(dst1Host);
172+ aclrtDestroyStream(stream);
173+ aclrtResetDevice(deviceId);
174+ aclFinalize();
175+ return 0;
176+}
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/data_reorder/data_utils.h+101-0
@@ -0,0 +1,101 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file data_utils.h
13+ * \brief
14+ */
15+ 
16+#ifndef DATA_UTILS_H
17+#define DATA_UTILS_H
18+#include <fcntl.h>
19+#include <sys/stat.h>
20+#include <unistd.h>
21+#include <fstream>
22+ 
23+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
24+ 
25+bool ReadFile(const std::string& filePath, size_t& fileSize, void* buffer, size_t bufferSize)
26+{
27+ struct stat sBuf;
28+ int fileStatus = stat(filePath.data(), &sBuf);
29+ if (fileStatus == -1) {
30+ ERROR_LOG("failed to get file");
31+ return false;
32+ }
33+ if (S_ISREG(sBuf.st_mode) == 0) {
34+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
35+ return false;
36+ }
37+ 
38+ std::ifstream file;
39+ file.open(filePath, std::ios::binary);
40+ if (!file.is_open()) {
41+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
42+ return false;
43+ }
44+ 
45+ std::filebuf* buf = file.rdbuf();
46+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
47+ if (size == 0) {
48+ ERROR_LOG("file size is 0");
49+ file.close();
50+ return false;
51+ }
52+ if (size > bufferSize) {
53+ ERROR_LOG("file size is larger than buffer size");
54+ file.close();
55+ return false;
56+ }
57+ if (buf->pubseekpos(0, std::ios::in) == std::streampos(std::streamoff(-1))) {
58+ ERROR_LOG("failed to seek file");
59+ file.close();
60+ return false;
61+ }
62+ if (buf->sgetn(static_cast<char*>(buffer), size) != static_cast<std::streamsize>(size)) {
63+ ERROR_LOG("failed to read file");
64+ file.close();
65+ return false;
66+ }
67+ fileSize = size;
68+ file.close();
69+ return true;
70+}
71+ 
72+/**
73+ * @brief Write data to file
74+ * @param [in] filePath: file path
75+ * @param [in] buffer: data to write to file
76+ * @param [in] size: size to write
77+ * @return write result
78+ */
79+bool WriteFile(const std::string& filePath, const void* buffer, size_t size)
80+{
81+ if (buffer == nullptr) {
82+ ERROR_LOG("Write file failed. buffer is nullptr");
83+ return false;
84+ }
85+ 
86+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
87+ if (fd < 0) {
88+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
89+ return false;
90+ }
91+ 
92+ size_t writeSize = write(fd, buffer, size);
93+ (void)close(fd);
94+ if (writeSize != size) {
95+ ERROR_LOG("Write file Failed.");
96+ return false;
97+ }
98+ 
99+ return true;
100+}
101+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/data_reorder/scripts/gen_data.py+45-0
@@ -0,0 +1,45 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+ 
19+def gen_golden_data_simple():
20+ total_length = 256
21+ data_type = np.float32
22+ src0 = np.arange(total_length, dtype=data_type)
23+ src1 = np.arange(1000, 1000 + total_length, dtype=data_type)
24+ combined = np.array([], dtype=data_type)
25+ for i in range(4):
26+ offset = i * 64
27+ tmp_src0 = src0[offset : offset + 64]
28+ tmp_src1 = src1[offset : offset + 64]
29+ combined = np.concatenate((combined, tmp_src0, tmp_src1))
30+ dst0 = combined[0::2]
31+ dst1 = combined[1::2]
32+ 
33+ os.makedirs("input", exist_ok=True)
34+ os.makedirs("output", exist_ok=True)
35+ src0.tofile("input/input_src0.bin")
36+ src1.tofile("input/input_src1.bin")
37+ dst0.tofile("output/golden_dst0.bin")
38+ dst1.tofile("output/golden_dst1.bin")
39+ 
40+ print("dst0:\n", dst0)
41+ print("dst1:\n", dst1)
42+ 
43+ 
44+if __name__ == "__main__":
45+ gen_golden_data_simple()
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/fill/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# ----------------------------------------------------------------------------------------------------------
6+cmake_minimum_required(VERSION 3.16)
7+ 
8+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, sim")
9+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
10+option(CANN_ASC_USE_EXPERIMENTAL "Enable experimental CANN ASC APIs" OFF)
11+ 
12+find_package(ASC REQUIRED)
13+ 
14+project(kernel_samples LANGUAGES ASC CXX)
15+ 
16+add_executable(demo
17+ fill.asc
18+)
19+ 
20+target_link_libraries(demo PRIVATE
21+)
22+ 
23+target_compile_options(demo PRIVATE
24+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
25+ $<$<AND:$<COMPILE_LANGUAGE:ASC>,$<BOOL:${CANN_ASC_USE_EXPERIMENTAL}>>:-DCANN_ASC_USE_EXPERIMENTAL>
26+)
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/fill/README.md+81-0
@@ -0,0 +1,81 @@
1+# fill样例
2+ 
3+## 概述
4+ 
5+本样例基于Tensor API实现fill运算（标量填充模式），主要调用实验性接口 `asc::te::experimental::fill`。
6+ 
7+## 本样例支持的产品及CANN软件版本
8+ 
9+| 产品 | CANN软件版本 |
10+|------|-------------|
11+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
12+ 
13+## 目录结构介绍
14+ 
15+```
16+├── fill
17+│ ├── scripts
18+│ │ │ ├── gen_data.py // 真值数据生成脚本
19+│ ├── CMakeLists.txt // 编译工程文件
20+│ ├── data_utils.h // 数据读入写出函数
21+│ ├── fill.asc // AscendC样例实现 & 调用样例
22+│ └── README.md // 样例介绍
23+```
24+ 
25+## 样例描述
26+- 样例功能：
27+ 将标量值1填充到输出向量的每个位置，向量元素个数为256，数据类型为float。
28+- 样例规格：
29+ <table>
30+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
31+ <tr><td rowspan="2" align="center">样例输出</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
32+ <tr><td align="center">y</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
33+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">fill</td></tr>
34+ </table>
35+- 样例实现：
36+ fill_vf函数内调用fill接口进行标量填充：
37+ - 使用fill接口将标量值填充到目的reg_tensor的每个元素
38+ - 使用store将结果写回UB
39+ - 调用实现
40+ 使用内核调用符<<<>>>调用核函数。
41+ 
42+## 编译运行
43+在本样例根目录下执行如下步骤，编译并执行样例。
44+- 配置环境变量
45+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/zh/quick_start.md#prepare&install)，配置环境变量，**当前仅支持使用[CANN master](../../../../../../docs/zh/quick_start.md#cann-install)**。
46+ 
47+ > **说明：** `${install_path}` 为CANN包安装目录，未指定安装目录时默认安装至 `/usr/local/Ascend` 下。
48+ 
49+- 样例执行
50+ 
51+ 在本样例目录下执行如下命令。
52+ ```bash
53+ mkdir -p build && cd build; # 创建并进入build目录
54+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # 编译工程（默认npu模式）
55+ python3 ../scripts/gen_data.py # 生成测试真值数据
56+ ./demo # 执行编译生成的可执行程序，执行样例
57+ ```
58+ 
59+ 使用 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
60+ 
61+ 示例如下：
62+ ```bash
63+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # NPU仿真模式
64+ ```
65+ 
66+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
67+ 
68+- 编译选项说明
69+ 
70+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
71+| ---------------------------| -----------------------------| ---------------------------------------------------|
72+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`sim` | 运行模式：NPU 运行、NPU仿真　　　　　　　|
73+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
74+| `CANN_ASC_USE_EXPERIMENTAL` | `ON`（本样例必选）、`OFF`（默认） | 开启实验性ASC接口 |
75+ 
76+- 执行结果
77+ 
78+ 执行结果如下，说明精度对比成功。
79+ ```bash
80+ test pass!
81+ ```
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/fill/README_en.md+86-0
@@ -0,0 +1,86 @@
1+# fill Example
2+ 
3+## Overview
4+ 
5+This example implements the `fill` operation in scalar fill mode based on the Tensor API. It primarily calls the experimental `asc::te::experimental::fill` interface.
6+ 
7+## Supported Products and CANN Versions
8+ 
9+| Product | CANN Version |
10+| --- | --- |
11+| Ascend 950PR/Ascend 950DT | >= CANN 9.1.0 |
12+ 
13+## Directory Structure
14+ 
15+```text
16+fill
17+├── scripts
18+│ └── gen_data.py // Script for generating ground truth data
19+├── CMakeLists.txt // Build project file
20+├── data_utils.h // Data read/write functions
21+├── fill.asc // Ascend C example implementation and invocation
22+└── README_en.md // Example documentation
23+```
24+ 
25+## Example Description
26+ 
27+- Function:
28+ Fills every position of the output vector with the scalar value 1. The vector contains 256 elements of the float data type.
29+- Specifications:
30+ <table>
31+ <tr><td rowspan="1" align="center">Example Type (OpType)</td><td colspan="3" align="center">AIV Example</td></tr>
32+ <tr><td rowspan="2" align="center">Output</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
33+ <tr><td align="center">y</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
34+ <tr><td rowspan="1" align="center">Kernel Function</td><td colspan="4" align="center">fill</td></tr>
35+ </table>
36+- Implementation:
37+ The `fill_vf` function calls the `fill` interface to perform scalar filling:
38+ - Use `fill` to write the scalar value to every element of the destination `reg_tensor`.
39+ - Use `store` to write the result back to UB.
40+ - Invoke the kernel function using the `<<<>>>` kernel launch syntax.
41+ 
42+## Build and Run
43+ 
44+Perform the following steps in the example root directory to build and run the example.
45+ 
46+- Configure environment variables
47+ Configure the environment variables according to the CANN development kit [installation instructions](../../../../../../docs/en/quick_start.md#prepare&install). **Currently, only [CANN master](../../../../../../docs/en/quick_start.md#cann-install) is supported.**
48+ 
49+ > **Note:** `${install_path}` is the CANN package installation directory. If no installation directory is specified, the default directory is `/usr/local/Ascend`.
50+ 
51+- Run the example
52+ 
53+ Run the following commands in the example directory:
54+ 
55+ ```bash
56+ mkdir -p build && cd build; # Create and enter the build directory
57+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # Build the project (NPU mode by default)
58+ python3 ../scripts/gen_data.py # Generate test ground truth data
59+ ./demo # Run the generated executable
60+ ```
61+ 
62+ To use NPU simulation mode, add the `-DCMAKE_ASC_RUN_MODE=sim` option.
63+ 
64+ Example:
65+ 
66+ ```bash
67+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DCANN_ASC_USE_EXPERIMENTAL=ON ..;make -j; # NPU simulation mode
68+ ```
69+ 
70+ > **Note:** Clear the CMake cache before switching build modes. Run `rm CMakeCache.txt` in the `build` directory, and then run CMake again.
71+ 
72+- Build options
73+ 
74+| Option | Values | Description |
75+| --- | --- | --- |
76+| `CMAKE_ASC_RUN_MODE` | `npu` (default), `sim` | Run mode: NPU execution or NPU simulation |
77+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU architecture: `dav-3510` corresponds to Ascend 950PR/Ascend 950DT |
78+| `CANN_ASC_USE_EXPERIMENTAL` | `ON` (required for this example), `OFF` (default) | Enables experimental ASC interfaces |
79+ 
80+- Execution result
81+ 
82+ The following output indicates that the accuracy comparison is successful:
83+ 
84+ ```bash
85+ test pass!
86+ ```
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/fill/data_utils.h+101-0
@@ -0,0 +1,101 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file data_utils.h
13+ * \brief
14+ */
15+ 
16+#ifndef DATA_UTILS_H
17+#define DATA_UTILS_H
18+#include <fcntl.h>
19+#include <sys/stat.h>
20+#include <unistd.h>
21+#include <fstream>
22+ 
23+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
24+ 
25+bool ReadFile(const std::string& filePath, size_t& fileSize, void* buffer, size_t bufferSize)
26+{
27+ struct stat sBuf;
28+ int fileStatus = stat(filePath.data(), &sBuf);
29+ if (fileStatus == -1) {
30+ ERROR_LOG("failed to get file");
31+ return false;
32+ }
33+ if (S_ISREG(sBuf.st_mode) == 0) {
34+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
35+ return false;
36+ }
37+ 
38+ std::ifstream file;
39+ file.open(filePath, std::ios::binary);
40+ if (!file.is_open()) {
41+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
42+ return false;
43+ }
44+ 
45+ std::filebuf* buf = file.rdbuf();
46+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
47+ if (size == 0) {
48+ ERROR_LOG("file size is 0");
49+ file.close();
50+ return false;
51+ }
52+ if (size > bufferSize) {
53+ ERROR_LOG("file size is larger than buffer size");
54+ file.close();
55+ return false;
56+ }
57+ if (buf->pubseekpos(0, std::ios::in) == std::streampos(std::streamoff(-1))) {
58+ ERROR_LOG("failed to seek file");
59+ file.close();
60+ return false;
61+ }
62+ if (buf->sgetn(static_cast<char*>(buffer), size) != static_cast<std::streamsize>(size)) {
63+ ERROR_LOG("failed to read file");
64+ file.close();
65+ return false;
66+ }
67+ fileSize = size;
68+ file.close();
69+ return true;
70+}
71+ 
72+/**
73+ * @brief Write data to file
74+ * @param [in] filePath: file path
75+ * @param [in] buffer: data to write to file
76+ * @param [in] size: size to write
77+ * @return write result
78+ */
79+bool WriteFile(const std::string& filePath, const void* buffer, size_t size)
80+{
81+ if (buffer == nullptr) {
82+ ERROR_LOG("Write file failed. buffer is nullptr");
83+ return false;
84+ }
85+ 
86+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
87+ if (fd < 0) {
88+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
89+ return false;
90+ }
91+ 
92+ size_t writeSize = write(fd, buffer, size);
93+ (void)close(fd);
94+ if (writeSize != size) {
95+ ERROR_LOG("Write file Failed.");
96+ return false;
97+ }
98+ 
99+ return true;
100+}
101+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/fill/fill.asc+132-0
@@ -0,0 +1,132 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <cmath>
12+#include <cstdint>
13+#include <cstdio>
14+#include <vector>
15+ 
16+#include "acl/acl.h"
17+#include "data_utils.h"
18+#include "tensor_api/experimental/vector_compute.h"
19+#include "tensor_api/tensor.h"
20+ 
21+constexpr uint32_t total_length = 256;
22+constexpr uint8_t buffer_lock_id = 0;
23+ 
24+template <typename T, typename U>
25+__simd_vf__ inline void fill_vf(T y_tensor, U scalar, uint32_t repeat_times, uint32_t one_repeat_size)
26+{
27+ uint32_t count = total_length;
28+ using data_type = typename T::data_type;
29+ for (uint16_t i = 0; i < repeat_times; ++i) {
30+ auto mask = asc::te::experimental::update_mask<data_type>(count);
31+ uint32_t offset = i * one_repeat_size;
32+ const auto coord = asc::te::make_coord(offset);
33+ auto y_reg = asc::te::experimental::fill(scalar, mask);
34+ asc::te::experimental::store(y_tensor, coord, y_reg);
35+ }
36+}
37+ 
38+template <typename T>
39+class Kernel {
40+public:
41+ __aicore__ inline void process(__gm__ T* y)
42+ {
43+ auto layout = asc::te::make_layout(asc::te::make_shape(total_length), asc::te::make_stride(1));
44+ auto y_gm = asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::gm>(y), layout);
45+ constexpr uint32_t y_offset = 0;
46+ auto y_ub = asc::te::make_tensor(asc::te::make_mem_ptr<asc::te::location::ub, T>(y_offset), layout);
47+ T scalar = 1;
48+ constexpr uint32_t one_repeat_size = asc_get_vf_len() / sizeof(T);
49+ uint16_t repeat_times =
50+ static_cast<uint16_t>(AscendC::Std::ceil_align(total_length, one_repeat_size) / one_repeat_size);
51+ asc_lock(pipe_t::PIPE_V, buffer_lock_id);
52+ fill_vf(y_ub, scalar, repeat_times, one_repeat_size);
53+ asc_unlock(pipe_t::PIPE_V, buffer_lock_id);
54+ asc_lock(pipe_t::PIPE_MTE3, buffer_lock_id);
55+ asc::te::copy(asc::te::make_copy(asc::te::copy_ub_to_gm{}), y_gm, y_ub);
56+ asc_unlock(pipe_t::PIPE_MTE3, buffer_lock_id);
57+ }
58+};
59+ 
60+template <typename T>
61+__global__ __vector__ void fill(__gm__ uint8_t* y)
62+{
63+ Kernel<T> kernel;
64+ kernel.process(reinterpret_cast<__gm__ T*>(y));
65+}
66+ 
67+static bool CompareResult(const void* outputData, uint32_t outSize)
68+{
69+ void* goldenData;
70+ aclrtMallocHost((void**)(&goldenData), outSize);
71+ size_t goldenSize = outSize;
72+ bool ret = ReadFile("./output/golden.bin", goldenSize, goldenData, goldenSize);
73+ if (ret) {
74+ printf("ReadFile golden.bin success!\n");
75+ } else {
76+ printf("test failed!\n");
77+ return false;
78+ }
79+ constexpr float eps = 1e-4;
80+ int64_t wrongNum = 0;
81+ 
82+ for (size_t i = 0; i < outSize / sizeof(float); i++) {
83+ float a = (reinterpret_cast<const float*>(outputData))[i];
84+ float b = (reinterpret_cast<const float*>(goldenData))[i];
85+ float ae = std::abs(a - b);
86+ float re = ae / std::abs(b);
87+ if (ae > eps && re > eps) {
88+ printf("CompareResult golden.bin failed at index %zu: output is %lf, golden is %lf\n", i, a, b);
89+ wrongNum++;
90+ }
91+ }
92+ aclrtFreeHost(goldenData);
93+ if (wrongNum != 0) {
94+ return false;
95+ } else {
96+ printf("CompareResult golden.bin success!\n");
97+ return true;
98+ }
99+}
100+ 
101+int32_t main(int32_t argc, char* argv[])
102+{
103+ uint32_t numBlocks = 1;
104+ aclInit(nullptr);
105+ int32_t deviceId = 0;
106+ aclrtSetDevice(deviceId);
107+ aclrtStream stream = nullptr;
108+ aclrtCreateStream(&stream);
109+ uint32_t totalLength = total_length;
110+ size_t outputByteSize = static_cast<size_t>(1) * totalLength * sizeof(float);
111+ uint8_t* yHost;
112+ uint8_t* yDevice;
113+ aclrtMallocHost((void**)(&yHost), outputByteSize);
114+ aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
115+ fill<float><<<numBlocks, 0, stream>>>(yDevice);
116+ aclrtSynchronizeStream(stream);
117+ aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
118+ WriteFile("./output/output.bin", yHost, outputByteSize);
119+ bool goldenResult = true;
120+ goldenResult = CompareResult(yHost, outputByteSize);
121+ if (goldenResult) {
122+ printf("test pass!\n");
123+ } else {
124+ printf("test failed!\n");
125+ }
126+ aclrtFree(yDevice);
127+ aclrtFreeHost(yHost);
128+ aclrtDestroyStream(stream);
129+ aclrtResetDevice(deviceId);
130+ aclFinalize();
131+ return 0;
132+}
Aexamples/01_simd_cpp_api/07_tensor_api/experimental/reg_vector_compute/fill/scripts/gen_data.py+29-0
@@ -0,0 +1,29 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+ 
19+def gen_golden_data_simple():
20+ total_length = 256
21+ data_type = np.float32
22+ fill_value = 1
23+ golden = np.full([1, total_length], fill_value, dtype=data_type)
24+ os.makedirs("output", exist_ok=True)
25+ golden.tofile("./output/golden.bin")
26+ 
27+ 
28+if __name__ == "__main__":
29+ gen_golden_data_simple()
Aimpl/tensor_api/experimental/arch/vector/compare_and_select_impl.h+56-0
@@ -0,0 +1,56 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#if !defined(ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS)
12+#warning \
13+ "impl/tensor_api/experimental/arch/vector/compare_and_select_impl.h is an internal header file and must not be used directly. Functions or variables defined in this file may be removed in the future."
14+#define ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
15+#define UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_COMPARE_AND_SELECT_IMPL_H
16+#endif
17+ 
18+#ifndef IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_COMPARE_AND_SELECT_IMPL_H
19+#define IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_COMPARE_AND_SELECT_IMPL_H
20+ 
21+#include <cstdint>
22+ 
23+namespace asc {
24+namespace te {
25+namespace experimental {
26+namespace detail {
27+template <typename T>
28+inline constexpr bool is_select_support_type = Std::is_one_of_v<
29+ T, bool, uint8_t, int8_t, fp8_e4m3fn_t, hifloat8_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, uint32_t,
30+ int32_t, float>;
31+}
32+ 
33+template <typename T>
34+__simd_callee__ inline reg_tensor<T> select(
35+ const reg_tensor<bool>& condition, const reg_tensor<T>& src0, const reg_tensor<T>& src1)
36+{
37+ static_assert(
38+ detail::is_select_support_type<T>,
39+ "select only supports bool, uint8_t, int8_t, "
40+ "fp8_e4m3fn_t, hifloat8_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t "
41+ "and float.");
42+ reg_tensor<T> dst;
43+ asc_select(dst.reg, src0.reg, src1.reg, condition.reg);
44+ dst.mask = src0.mask;
45+ return dst;
46+}
47+} // namespace experimental
48+} // namespace te
49+} // namespace asc
50+ 
51+#endif // IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_COMPARE_AND_SELECT_IMPL_H
52+ 
53+#if defined(UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_COMPARE_AND_SELECT_IMPL_H)
54+#undef ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
55+#undef UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_COMPARE_AND_SELECT_IMPL_H
56+#endif
Aimpl/tensor_api/experimental/arch/vector/data_padding_impl.h+71-0
@@ -0,0 +1,71 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#if !defined(ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS)
12+#warning \
13+ "impl/tensor_api/experimental/arch/vector/data_padding_impl.h is an internal header file and must not be used directly. Functions or variables defined in this file may be removed in the future."
14+#define ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
15+#define UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_PADDING_IMPL_H
16+#endif
17+ 
18+#ifndef IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_PADDING_IMPL_H
19+#define IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_PADDING_IMPL_H
20+ 
21+#include <cstdint>
22+ 
23+namespace asc {
24+namespace te {
25+namespace experimental {
26+namespace detail {
27+template <typename T>
28+inline constexpr bool is_fill_support_type = Std::is_one_of_v<
29+ T, uint8_t, int8_t, fp8_e4m3fn_t, fp8_e8m0_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t,
30+ float>;
31+}
32+ 
33+template <size_t N, typename T>
34+__simd_callee__ inline reg_tensor<T> fill(const T& value, const reg_tensor<bool>& mask)
35+{
36+ static_assert(
37+ detail::is_fill_support_type<T>,
38+ "fill only supports uint8_t, int8_t, "
39+ "fp8_e4m3fn_t, fp8_e8m0_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t "
40+ "and float.");
41+ static_assert(N == 1, "fill only supports N==1.");
42+ reg_tensor<T> dst;
43+ asc_duplicate_scalar(dst.reg, value, mask.reg);
44+ dst.mask = mask.reg;
45+ return dst;
46+}
47+ 
48+template <size_t N, typename T>
49+__simd_callee__ inline reg_tensor<T> fill(const reg_tensor<T>& src, const reg_tensor<bool>& mask)
50+{
51+ static_assert(
52+ detail::is_fill_support_type<T>,
53+ "fill only supports uint8_t, int8_t, "
54+ "fp8_e4m3fn_t, fp8_e8m0_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t "
55+ "and float.");
56+ static_assert(N == 1, "fill only supports N==1.");
57+ reg_tensor<T> dst;
58+ asc_duplicate(dst.reg, src.reg, mask.reg);
59+ dst.mask = mask.reg;
60+ return dst;
61+}
62+} // namespace experimental
63+} // namespace te
64+} // namespace asc
65+ 
66+#endif // IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_PADDING_IMPL_H
67+ 
68+#if defined(UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_PADDING_IMPL_H)
69+#undef ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
70+#undef UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_PADDING_IMPL_H
71+#endif
Aimpl/tensor_api/experimental/arch/vector/data_reorder_impl.h+72-0
@@ -0,0 +1,72 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#if !defined(ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS)
12+#warning \
13+ "impl/tensor_api/experimental/arch/vector/data_reorder_impl.h is an internal header file and must not be used directly. Functions or variables defined in this file may be removed in the future. Please use tensor_api/experimental/arch/vector/data_reorder.h instead."
14+#define ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
15+#define UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_REORDER_IMPL_H
16+#endif
17+ 
18+#ifndef IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_REORDER_IMPL_H
19+#define IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_REORDER_IMPL_H
20+ 
21+#include "tensor_api/experimental/arch/vector/data_reorder.h"
22+#include "tensor_api/experimental/arch/vector/mask_reg_compute.h"
23+ 
24+namespace asc {
25+namespace te {
26+namespace experimental {
27+namespace detail {
28+template <typename T>
29+inline constexpr bool is_data_reorder_support_type = AscendC::Std::is_one_of_v<
30+ T, uint8_t, int8_t, fp8_e4m3fn_t, fp8_e8m0_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, uint32_t, int32_t,
31+ float>;
32+} // namespace detail
33+ 
34+template <typename T>
35+__simd_callee__ inline reg_pair<T> interleave(reg_tensor<T> src0, reg_tensor<T> src1)
36+{
37+ static_assert(
38+ detail::is_data_reorder_support_type<T>,
39+ "interleave with reg_tensor<T> only supports "
40+ "uint8_t, int8_t, fp8_e4m3fn_t, fp8_e8m0_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, "
41+ "uint32_t, int32_t and float.");
42+ reg_pair<T> dst;
43+ asc_intlv(dst.first.reg, dst.second.reg, src0.reg, src1.reg);
44+ dst.first.with_mask(all_mask<T>());
45+ dst.second.mask = dst.first.mask;
46+ return dst;
47+}
48+ 
49+template <typename T>
50+__simd_callee__ inline reg_pair<T> deinterleave(reg_tensor<T> src0, reg_tensor<T> src1)
51+{
52+ static_assert(
53+ detail::is_data_reorder_support_type<T>,
54+ "deinterleave with reg_tensor<T> only supports "
55+ "uint8_t, int8_t, fp8_e4m3fn_t, fp8_e8m0_t, fp8_e5m2_t, uint16_t, int16_t, half, bfloat16_t, "
56+ "uint32_t, int32_t and float.");
57+ reg_pair<T> dst;
58+ asc_deintlv(dst.first.reg, dst.second.reg, src0.reg, src1.reg);
59+ dst.first.with_mask(all_mask<T>());
60+ dst.second.mask = dst.first.mask;
61+ return dst;
62+}
63+} // namespace experimental
64+} // namespace te
65+} // namespace asc
66+ 
67+#endif // IMPL_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_REORDER_IMPL_H
68+ 
69+#if defined(UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_REORDER_IMPL_H)
70+#undef ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
71+#undef UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_REORDER_IMPL_H
72+#endif
Ainclude/tensor_api/experimental/arch/vector/compare_and_select.h+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#if !defined(ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS)
12+#define ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
13+#define UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_COMPARE_AND_SELECT_H
14+#endif
15+ 
16+#ifndef INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_COMPARE_AND_SELECT_H
17+#define INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_COMPARE_AND_SELECT_H
18+ 
19+#include "tensor_api/experimental/arch/vector/reg_tensor.h"
20+ 
21+namespace asc {
22+namespace te {
23+namespace experimental {
24+/**
25+ * @brief Selects elements from two register tensors according to a condition tensor.
26+ * @param condition Per-element selection condition.
27+ * @param src0 Source tensor selected when the corresponding condition is true.
28+ * @param src1 Source tensor selected when the corresponding condition is false.
29+ * @return A register tensor containing the selected elements.
30+ */
31+template <typename T>
32+__simd_callee__ inline reg_tensor<T> select(
33+ const reg_tensor<bool>& condition, const reg_tensor<T>& src0, const reg_tensor<T>& src1);
34+} // namespace experimental
35+} // namespace te
36+} // namespace asc
37+ 
38+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
39+#include "impl/tensor_api/experimental/arch/vector/compare_and_select_impl.h"
40+#endif
41+ 
42+#endif // INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_COMPARE_AND_SELECT_H
43+ 
44+#if defined(UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_COMPARE_AND_SELECT_H)
45+#undef ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
46+#undef UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_COMPARE_AND_SELECT_H
47+#endif
Ainclude/tensor_api/experimental/arch/vector/data_padding.h+61-0
@@ -0,0 +1,61 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#if !defined(ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS)
12+#define ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
13+#define UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_PADDING_H
14+#endif
15+ 
16+#ifndef INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_PADDING_H
17+#define INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_PADDING_H
18+ 
19+#include <cstddef>
20+ 
21+#include "tensor_api/experimental/arch/vector/reg_tensor.h"
22+#include "tensor_api/experimental/arch/vector/mask_reg_compute.h"
23+ 
24+namespace asc {
25+namespace te {
26+namespace experimental {
27+/**
28+ * @brief Fills masked elements of a register tensor with a scalar value.
29+ * @tparam N Number of register tensors to fill. Only 1 is supported.
30+ * @tparam T Element type of the register tensor.
31+ * @param value Scalar value copied to each enabled element.
32+ * @param mask Mask that specifies the elements to fill.
33+ * @return A register tensor containing the filled elements and the input mask.
34+ */
35+template <size_t N = 1, typename T>
36+__simd_callee__ inline reg_tensor<T> fill(const T& value, const reg_tensor<bool>& mask = all_mask<T>());
37+ 
38+/**
39+ * @brief Copies masked elements from a register tensor.
40+ * @tparam N Number of register tensors to fill. Only 1 is supported.
41+ * @tparam T Element type of the register tensor.
42+ * @param val Source register tensor.
43+ * @param mask Mask that specifies the elements to copy.
44+ * @return A register tensor containing the copied elements and the input mask.
45+ */
46+template <size_t N = 1, typename T>
47+__simd_callee__ inline reg_tensor<T> fill(const reg_tensor<T>& src, const reg_tensor<bool>& mask = all_mask<T>());
48+} // namespace experimental
49+} // namespace te
50+} // namespace asc
51+ 
52+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
53+#include "impl/tensor_api/experimental/arch/vector/data_padding_impl.h"
54+#endif
55+ 
56+#endif // INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_PADDING_H
57+ 
58+#if defined(UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_PADDING_H)
59+#undef ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
60+#undef UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_PADDING_H
61+#endif
Ainclude/tensor_api/experimental/arch/vector/data_reorder.h+56-0
@@ -0,0 +1,56 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#if !defined(ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS)
12+#define ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
13+#define UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_REORDER_H
14+#endif
15+ 
16+#ifndef INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_REORDER_H
17+#define INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_REORDER_H
18+ 
19+#include "tensor_api/experimental/arch/vector/reg_tensor.h"
20+ 
21+namespace asc {
22+namespace te {
23+namespace experimental {
24+/**
25+ * @brief Interleaves elements from two register tensors.
26+ * @tparam T Element type of the register tensors.
27+ * @param src0 First source register tensor.
28+ * @param src1 Second source register tensor.
29+ * @return A pair of register tensors containing the interleaved elements. Both result tensors use a full mask.
30+ */
31+template <typename T>
32+__simd_callee__ inline reg_pair<T> interleave(reg_tensor<T> src0, reg_tensor<T> src1);
33+ 
34+/**
35+ * @brief Deinterleaves elements from two register tensors.
36+ * @tparam T Element type of the register tensors.
37+ * @param src0 First source register tensor.
38+ * @param src1 Second source register tensor.
39+ * @return A pair of register tensors containing the deinterleaved elements. Both result tensors use a full mask.
40+ */
41+template <typename T>
42+__simd_callee__ inline reg_pair<T> deinterleave(reg_tensor<T> src0, reg_tensor<T> src1);
43+} // namespace experimental
44+} // namespace te
45+} // namespace asc
46+ 
47+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
48+#include "impl/tensor_api/experimental/arch/vector/data_reorder_impl.h"
49+#endif
50+ 
51+#endif // INCLUDE_TENSOR_API_EXPERIMENTAL_ARCH_VECTOR_DATA_REORDER_H
52+ 
53+#if defined(UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_REORDER_H)
54+#undef ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS
55+#undef UNDEF_ASCENDC_TENSOR_API_INCLUDE_COMPILER_INTERNAL_HEADERS_DATA_REORDER_H
56+#endif
Minclude/tensor_api/experimental/vector_compute.h+3-0
@@ -31,6 +31,9 @@
31#include "tensor_api/experimental/arch/vector/mask_reg_compute.h"31#include "tensor_api/experimental/arch/vector/mask_reg_compute.h"
32#include "tensor_api/experimental/arch/vector/reg_data_load.h"32#include "tensor_api/experimental/arch/vector/reg_data_load.h"
33#include "tensor_api/experimental/arch/vector/reg_data_store.h"33#include "tensor_api/experimental/arch/vector/reg_data_store.h"
34+#include "tensor_api/experimental/arch/vector/data_reorder.h"
35+#include "tensor_api/experimental/arch/vector/data_padding.h"
36+#include "tensor_api/experimental/arch/vector/compare_and_select.h"
34 37 
35#endif // INCLUDE_TENSOR_API_EXPERIMENTAL_VECTOR_COMPUTE_H38#endif // INCLUDE_TENSOR_API_EXPERIMENTAL_VECTOR_COMPUTE_H
36 39 
Atests/api/tensor_api/npu_arch_3510/aiv/test_tensor_api_compare_and_select.cpp+53-0
@@ -0,0 +1,53 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+ 
13+#include "tensor_api/stub/cce_stub.h"
14+#include "tensor_api/experimental/vector_compute.h"
15+#include "tensor_api/tensor.h"
16+ 
17+class TensorApiVectorCompareAndSelect3510 : public testing::Test {};
18+ 
19+template <typename T>
20+__aicore__ inline void TestSelect()
21+{
22+ asc::te::experimental::reg_tensor<bool> condition{};
23+ asc::te::experimental::reg_tensor<T> src0{};
24+ asc::te::experimental::reg_tensor<T> src1{};
25+ src0.mask = asc::te::experimental::all_mask<T>().reg;
26+ 
27+ auto dst = asc::te::experimental::select(condition, src0, src1);
28+ static_assert(AscendC::Std::is_same_v<decltype(dst), asc::te::experimental::reg_tensor<T>>);
29+ EXPECT_EQ(dst.mask, src0.mask);
30+}
31+ 
32+#define COMPARE_AND_SELECT_TEST(Function, DataType) \
33+ TEST_F(TensorApiVectorCompareAndSelect3510, Function##_##DataType) \
34+ { \
35+ Test##Function<DataType>(); \
36+ SUCCEED(); \
37+ }
38+ 
39+#define SELECT_TYPE_TEST(DataType) COMPARE_AND_SELECT_TEST(Select, DataType)
40+ 
41+SELECT_TYPE_TEST(bool)
42+SELECT_TYPE_TEST(uint8_t)
43+SELECT_TYPE_TEST(int8_t)
44+SELECT_TYPE_TEST(fp8_e4m3fn_t)
45+SELECT_TYPE_TEST(hifloat8_t)
46+SELECT_TYPE_TEST(fp8_e5m2_t)
47+SELECT_TYPE_TEST(uint16_t)
48+SELECT_TYPE_TEST(int16_t)
49+SELECT_TYPE_TEST(half)
50+SELECT_TYPE_TEST(bfloat16_t)
51+SELECT_TYPE_TEST(uint32_t)
52+SELECT_TYPE_TEST(int32_t)
53+SELECT_TYPE_TEST(float)
Atests/api/tensor_api/npu_arch_3510/aiv/test_tensor_api_data_padding.cpp+75-0
@@ -0,0 +1,75 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+ 
13+#include "tensor_api/stub/cce_stub.h"
14+#include "tensor_api/experimental/vector_compute.h"
15+#include "tensor_api/tensor.h"
16+ 
17+class TensorApiVectorDataPadding3510 : public testing::Test {};
18+ 
19+template <typename T>
20+__aicore__ inline void TestFillScalar()
21+{
22+ auto mask = asc::te::experimental::make_mask<asc::te::experimental::mask_pattern::every3, T>();
23+ auto dst = asc::te::experimental::fill(T{}, mask);
24+ static_assert(AscendC::Std::is_same_v<decltype(dst), asc::te::experimental::reg_tensor<T>>);
25+ EXPECT_EQ(dst.mask, mask.reg);
26+}
27+ 
28+template <typename T>
29+__aicore__ inline void TestFillRegTensor()
30+{
31+ asc::te::experimental::reg_tensor<T> src{};
32+ auto mask = asc::te::experimental::make_mask<asc::te::experimental::mask_pattern::every3, T>();
33+ auto dst = asc::te::experimental::fill(src, mask);
34+ static_assert(AscendC::Std::is_same_v<decltype(dst), asc::te::experimental::reg_tensor<T>>);
35+ EXPECT_EQ(dst.mask, mask.reg);
36+}
37+ 
38+template <typename T>
39+__aicore__ inline void TestFillDefaultMask()
40+{
41+ asc::te::experimental::reg_tensor<T> src{};
42+ 
43+ auto scalar_dst = asc::te::experimental::fill(T{});
44+ auto reg_tensor_dst = asc::te::experimental::fill(src);
45+ static_assert(AscendC::Std::is_same_v<decltype(scalar_dst), asc::te::experimental::reg_tensor<T>>);
46+ static_assert(AscendC::Std::is_same_v<decltype(reg_tensor_dst), asc::te::experimental::reg_tensor<T>>);
47+ auto mask = asc::te::experimental::all_mask<T>();
48+ EXPECT_EQ(scalar_dst.mask, mask.reg);
49+ EXPECT_EQ(reg_tensor_dst.mask, mask.reg);
50+}
51+ 
52+#define DATA_PADDING_TEST(Function, DataType) \
53+ TEST_F(TensorApiVectorDataPadding3510, Function##_##DataType) \
54+ { \
55+ Test##Function<DataType>(); \
56+ SUCCEED(); \
57+ }
58+ 
59+#define DATA_PADDING_TYPE_TESTS(DataType) \
60+ DATA_PADDING_TEST(FillScalar, DataType) \
61+ DATA_PADDING_TEST(FillRegTensor, DataType) \
62+ DATA_PADDING_TEST(FillDefaultMask, DataType)
63+ 
64+DATA_PADDING_TYPE_TESTS(uint8_t)
65+DATA_PADDING_TYPE_TESTS(int8_t)
66+DATA_PADDING_TYPE_TESTS(fp8_e4m3fn_t)
67+DATA_PADDING_TYPE_TESTS(fp8_e8m0_t)
68+DATA_PADDING_TYPE_TESTS(fp8_e5m2_t)
69+DATA_PADDING_TYPE_TESTS(uint16_t)
70+DATA_PADDING_TYPE_TESTS(int16_t)
71+DATA_PADDING_TYPE_TESTS(half)
72+DATA_PADDING_TYPE_TESTS(bfloat16_t)
73+DATA_PADDING_TYPE_TESTS(uint32_t)
74+DATA_PADDING_TYPE_TESTS(int32_t)
75+DATA_PADDING_TYPE_TESTS(float)
Atests/api/tensor_api/npu_arch_3510/aiv/test_tensor_api_data_reorder.cpp+65-0
@@ -0,0 +1,65 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+ 
13+#include "tensor_api/stub/cce_stub.h"
14+#include "tensor_api/experimental/vector_compute.h"
15+#include "tensor_api/tensor.h"
16+ 
17+class TensorApiVectorDataReorder3510 : public testing::Test {};
18+ 
19+template <typename T>
20+__aicore__ inline void TestInterleave()
21+{
22+ asc::te::experimental::reg_tensor<T> src0{};
23+ asc::te::experimental::reg_tensor<T> src1{};
24+ 
25+ auto [dst0, dst1] = asc::te::experimental::interleave<T>(src0, src1);
26+ auto mask = asc::te::experimental::all_mask<T>();
27+ EXPECT_EQ(dst0.mask, mask.reg);
28+ EXPECT_EQ(dst1.mask, mask.reg);
29+}
30+ 
31+template <typename T>
32+__aicore__ inline void TestDeinterleave()
33+{
34+ asc::te::experimental::reg_tensor<T> src0{};
35+ asc::te::experimental::reg_tensor<T> src1{};
36+ 
37+ auto mask = asc::te::experimental::all_mask<T>();
38+ auto [dst0, dst1] = asc::te::experimental::deinterleave<T>(src0, src1);
39+ EXPECT_EQ(dst0.mask, mask.reg);
40+ EXPECT_EQ(dst1.mask, mask.reg);
41+}
42+ 
43+#define DATA_REORDER_TEST(Function, DataType) \
44+ TEST_F(TensorApiVectorDataReorder3510, Function##_##DataType) \
45+ { \
46+ Test##Function<DataType>(); \
47+ SUCCEED(); \
48+ }
49+ 
50+#define DATA_REORDER_TYPE_TESTS(DataType) \
51+ DATA_REORDER_TEST(Interleave, DataType) \
52+ DATA_REORDER_TEST(Deinterleave, DataType)
53+ 
54+DATA_REORDER_TYPE_TESTS(uint8_t)
55+DATA_REORDER_TYPE_TESTS(int8_t)
56+DATA_REORDER_TYPE_TESTS(fp8_e4m3fn_t)
57+DATA_REORDER_TYPE_TESTS(fp8_e8m0_t)
58+DATA_REORDER_TYPE_TESTS(fp8_e5m2_t)
59+DATA_REORDER_TYPE_TESTS(uint16_t)
60+DATA_REORDER_TYPE_TESTS(int16_t)
61+DATA_REORDER_TYPE_TESTS(half)
62+DATA_REORDER_TYPE_TESTS(bfloat16_t)
63+DATA_REORDER_TYPE_TESTS(uint32_t)
64+DATA_REORDER_TYPE_TESTS(int32_t)
65+DATA_REORDER_TYPE_TESTS(float)