Star704
1007
代码
代码
Issues82
Pull Requests181
Actions
项目讨论
Wiki
分析
Star704
1007
已合并
feat: 新增C API Load类指令返回值接口 #5307
lihuaichao创建于 6 天前
feat: 新增C API Load类指令返回值接口 #5307
已合并
lihuaichao创建于 6 天前
从lihuaichao/asc-devkit_950capi_load_return合入到cann/asc-devkitmaster
共 22 个文件变更+2539-57
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_load.md+13-4
@@ -26,25 +26,34 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中按dtype对齐的起始地址读取VL长度数据，并搬入矢量数据寄存器，搬运过程中数据格式和内容保持不变。连续搬入时，需要在每次调用前手动更新源地址。该接口为易用性接口，对性能有要求时可使用[asc_loadunalign](asc_loadunalign.md)或[asc_loadunalign_postupdate](asc_loadunalign_postupdate.md)。29+从Unified Buffer（UB）中按dtype对齐的起始地址读取VL长度数据，并通过函数返回值返回或写入目的矢量数据寄存器，搬运过程中数据格式和内容保持不变。连续搬入时，需要在每次调用前手动更新源地址。该接口为易用性接口，对性能有要求时可使用[asc_loadunalign](asc_loadunalign.md)或[asc_loadunalign_postupdate](asc_loadunalign_postupdate.md)。
30 30 
31本接口仅在AIV上生效，非AIV调用直接返回。31本接口仅在AIV上生效，非AIV调用直接返回。
32 32 
33## 函数原型33## 函数原型
34 34 
35```c35```c
36+// 通过函数返回值返回结果。
37+__simd_callee__ inline vector_<dtype> asc_load(__ubuf__ <dtype>* src)
38+ 
39+// 通过引用参数输出结果。
36__simd_callee__ inline void asc_load(vector_<dtype>& dst,40__simd_callee__ inline void asc_load(vector_<dtype>& dst,
37 __ubuf__ <dtype>* src)41 __ubuf__ <dtype>* src)
38```42```
39 43 
40### dtype支持数据类型44### dtype支持数据类型
41 45 
42-dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`、`int64_t`。当dtype为`int4b_t`时，dst的实际类型为`vector_int4x2_t`。46+- 返回值类型接口支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。
47+- 无返回值类型接口支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`、`int64_t`。
48+ 
49+当dtype为`int4b_t`时，目的矢量数据寄存器的实际类型为`vector_int4x2_t`。
43 50 
44### 函数原型典型示例51### 函数原型典型示例
45 52 
46```c53```c
47// 示例：float类型。54// 示例：float类型。
55+__simd_callee__ inline vector_float asc_load(__ubuf__ float* src)
56+ 
48__simd_callee__ inline void asc_load(vector_float& dst,57__simd_callee__ inline void asc_load(vector_float& dst,
49 __ubuf__ float* src)58 __ubuf__ float* src)
50```59```
@@ -55,14 +64,14 @@ __simd_callee__ inline void asc_load(vector_float& dst,
55 64 
56| 参数名 | 输入/输出 | 描述 |65| 参数名 | 输入/输出 | 描述 |
57|---|---|---|66|---|---|---|
58-| dst | 输出 | 目的矢量数据寄存器。dtype必须与`src`一致，搬入VL长度数据。 |67+| dst | 输出 | 目的矢量数据寄存器。仅无返回值类型接口包含该参数。dtype必须与`src`一致，搬入VL长度数据。 |
59| src | 输入 | 源UB地址，实际读取地址必须按dtype对齐。 |68| src | 输入 | 源UB地址，实际读取地址必须按dtype对齐。 |
60 69 
61矢量数据寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。70矢量数据寄存器的详细说明请参见[reg数据类型定义](../../defs/type/data_type_definition.md)。
62 71 
63## 返回值说明72## 返回值说明
64 73 
65-无74+对于返回值类型接口，返回保存搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
66 75 
67## 约束说明76## 约束说明
68 77 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign.md+14-4
@@ -26,13 +26,15 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中32字节对齐的起始地址读取数据，并搬入矢量数据寄存器或掩码寄存器，搬运过程中数据格式和内容保持不变。本接口提供四种功能模式：29+从Unified Buffer（UB）中32字节对齐的起始地址读取数据，并通过函数返回值返回或写入目的矢量数据寄存器或掩码寄存器，搬运过程中数据格式和内容保持不变。本接口提供四种功能模式：
30 30 
31- **连续对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器或掩码寄存器，由用户自行更新源地址。31- **连续对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器或掩码寄存器，由用户自行更新源地址。
32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
33- **地址寄存器偏移搬入模式**：通过地址寄存器指定相对源起始地址的偏移，常用于Hardware Loop内偏移随循环计数变化的对齐搬入场景。需要与[asc_update_addr_reg](../reg_addr_reg/asc_update_addr_reg.md)配合使用。33- **地址寄存器偏移搬入模式**：通过地址寄存器指定相对源起始地址的偏移，常用于Hardware Loop内偏移随循环计数变化的对齐搬入场景。需要与[asc_update_addr_reg](../reg_addr_reg/asc_update_addr_reg.md)配合使用。
34- **非连续对齐搬入模式**：单条指令非连续搬入8个`DataBlock`，每个`DataBlock`的数据量为32字节。支持配置数据块之间的地址步长和起始读取位置。34- **非连续对齐搬入模式**：单条指令非连续搬入8个`DataBlock`，每个`DataBlock`的数据量为32字节。支持配置数据块之间的地址步长和起始读取位置。
35 35 
36+连续对齐搬入模式中，通过函数返回值返回掩码寄存器时，请使用[asc_loadalign_mask](asc_loadalign_mask.md)。非连续对齐搬入模式中，通过函数返回值返回矢量数据寄存器时，请使用[asc_loadalign_datablock_strided](asc_loadalign_datablock_strided.md)。
37+ 
36本接口仅在AIV上生效，非AIV调用直接返回。38本接口仅在AIV上生效，非AIV调用直接返回。
37 39 
38## 函数原型40## 函数原型
@@ -40,6 +42,9 @@
40### 连续对齐搬入模式42### 连续对齐搬入模式
41 43 
42```c44```c
45+// 通过函数返回值返回矢量数据寄存器。
46+__simd_callee__ inline vector_<dtype> asc_loadalign(__ubuf__ <dtype>* src)
47+ 
43// 搬入矢量数据寄存器。48// 搬入矢量数据寄存器。
44__simd_callee__ inline void asc_loadalign(vector_<dtype>& dst,49__simd_callee__ inline void asc_loadalign(vector_<dtype>& dst,
45 __ubuf__ <dtype>* src)50 __ubuf__ <dtype>* src)
@@ -50,12 +55,17 @@ __simd_callee__ inline void asc_loadalign(vector_bool& dst,
50 55 
51#### dtype支持数据类型56#### dtype支持数据类型
52 57 
53-dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`、`int64_t`、`uint64_t`。当dtype为`int4b_t`时，dst的实际类型为`vector_int4x2_t`。58+- 返回值类型接口支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。
59+- 无返回值类型接口支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`、`int64_t`、`uint64_t`。
60+ 
61+当dtype为`int4b_t`时，目的矢量数据寄存器的实际类型为`vector_int4x2_t`。
54 62 
55#### 函数原型典型示例63#### 函数原型典型示例
56 64 
57```c65```c
58// 示例：float类型。66// 示例：float类型。
67+__simd_callee__ inline vector_float asc_loadalign(__ubuf__ float* src)
68+ 
59__simd_callee__ inline void asc_loadalign(vector_float& dst,69__simd_callee__ inline void asc_loadalign(vector_float& dst,
60 __ubuf__ float* src)70 __ubuf__ float* src)
61```71```
@@ -145,7 +155,7 @@ __simd_callee__ inline void asc_loadalign(vector_float& dst,
145 155 
146| 参数名 | 输入/输出 | 描述 |156| 参数名 | 输入/输出 | 描述 |
147|---|---|---|157|---|---|---|
148-| dst | 输出 | 目的矢量数据寄存器或掩码寄存器。<br>&bull; 当`dst`为矢量数据寄存器，`dtype`必须与`src`一致，搬入VL长度数据。<br>&bull; 当`dst`为掩码寄存器，搬入VL/8长度数据。 |158+| dst | 输出 | 目的矢量数据寄存器或掩码寄存器，仅无返回值类型接口包含该参数。<br>&bull; 当`dst`为矢量数据寄存器，`dtype`必须与`src`一致，搬入VL长度数据。<br>&bull; 当`dst`为掩码寄存器，搬入VL/8长度数据。 |
149| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |159| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |
150 160 
151### 立即数偏移搬入模式161### 立即数偏移搬入模式
@@ -184,7 +194,7 @@ __simd_callee__ inline void asc_loadalign(vector_float& dst,
184 194 
185## 返回值说明195## 返回值说明
186 196 
187-无197+对于返回值类型接口，返回保存连续对齐搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
188 198 
189## 约束说明199## 约束说明
190 200 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_datablock.md+9-3
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中32字节对齐的起始地址读取一个`DataBlock`（32字节），并将该`DataBlock`广播到整个矢量数据寄存器。搬运过程中数据格式和内容保持不变。本接口提供三种功能模式：29+从Unified Buffer（UB）中32字节对齐的起始地址读取一个`DataBlock`（32字节），并将该`DataBlock`广播到整个矢量数据寄存器。结果通过函数返回值返回或写入目的矢量数据寄存器，搬运过程中数据格式和内容保持不变。本接口提供三种功能模式：
30 30 
31- **对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器，由用户自行更新源地址。31- **对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器，由用户自行更新源地址。
32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
@@ -39,6 +39,10 @@
39### 对齐搬入模式39### 对齐搬入模式
40 40 
41```c41```c
42+// 通过函数返回值返回结果。
43+__simd_callee__ inline vector_<dtype> asc_loadalign_brc_datablock(__ubuf__ <dtype>* src)
44+ 
45+// 通过引用参数输出结果。
42__simd_callee__ inline void asc_loadalign_brc_datablock(vector_<dtype>& dst,46__simd_callee__ inline void asc_loadalign_brc_datablock(vector_<dtype>& dst,
43 __ubuf__ <dtype>* src)47 __ubuf__ <dtype>* src)
44```48```
@@ -51,6 +55,8 @@ dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`
51 55 
52```c56```c
53// 示例：half类型。57// 示例：half类型。
58+__simd_callee__ inline vector_half asc_loadalign_brc_datablock(__ubuf__ half* src)
59+ 
54__simd_callee__ inline void asc_loadalign_brc_datablock(vector_half& dst,60__simd_callee__ inline void asc_loadalign_brc_datablock(vector_half& dst,
55 __ubuf__ half* src)61 __ubuf__ half* src)
56```62```
@@ -105,7 +111,7 @@ __simd_callee__ inline void asc_loadalign_brc_datablock(vector_half& dst,
105 111 
106| 参数名 | 输入/输出 | 描述 |112| 参数名 | 输入/输出 | 描述 |
107|---|---|---|113|---|---|---|
108-| dst | 输出 | 目的矢量数据寄存器。dtype必须与`src`一致，搬入VL长度数据。 |114+| dst | 输出 | 目的矢量数据寄存器。仅无返回值类型接口包含该参数。dtype必须与`src`一致，搬入VL长度数据。 |
109| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |115| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |
110 116 
111### 立即数偏移搬入模式117### 立即数偏移搬入模式
@@ -132,7 +138,7 @@ __simd_callee__ inline void asc_loadalign_brc_datablock(vector_half& dst,
132 138 
133## 返回值说明139## 返回值说明
134 140 
135-无141+对于返回值类型接口，返回保存广播搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
136 142 
137## 约束说明143## 约束说明
138 144 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_elem.md+9-3
@@ -25,7 +25,7 @@
25<!-- end id7 -->25<!-- end id7 -->
26 26 
27## 功能说明27## 功能说明
28-从Unified Buffer（UB）中按dtype对齐的起始地址读取一个元素，并将该元素广播到整个矢量数据寄存器。搬运过程中数据格式和内容保持不变。本接口提供三种功能模式：28+从Unified Buffer（UB）中按dtype对齐的起始地址读取一个元素，并将该元素广播到整个矢量数据寄存器。结果通过函数返回值返回或写入目的矢量数据寄存器，搬运过程中数据格式和内容保持不变。本接口提供三种功能模式：
29 29 
30- **对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器，由用户自行更新源地址。30- **对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器，由用户自行更新源地址。
31- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。31- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
@@ -38,6 +38,10 @@
38### 对齐搬入模式38### 对齐搬入模式
39 39 
40```c40```c
41+// 通过函数返回值返回结果。
42+__simd_callee__ inline vector_<dtype> asc_loadalign_brc_elem(__ubuf__ <dtype>* src)
43+ 
44+// 通过引用参数输出结果。
41__simd_callee__ inline void asc_loadalign_brc_elem(vector_<dtype>& dst,45__simd_callee__ inline void asc_loadalign_brc_elem(vector_<dtype>& dst,
42 __ubuf__ <dtype>* src)46 __ubuf__ <dtype>* src)
43```47```
@@ -50,6 +54,8 @@ dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`
50 54 
51```c55```c
52// 示例：int8_t类型。56// 示例：int8_t类型。
57+__simd_callee__ inline vector_int8_t asc_loadalign_brc_elem(__ubuf__ int8_t* src)
58+ 
53__simd_callee__ inline void asc_loadalign_brc_elem(vector_int8_t& dst,59__simd_callee__ inline void asc_loadalign_brc_elem(vector_int8_t& dst,
54 __ubuf__ int8_t* src)60 __ubuf__ int8_t* src)
55```61```
@@ -104,7 +110,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem(vector_half& dst,
104 110 
105| 参数名 | 输入/输出 | 描述 |111| 参数名 | 输入/输出 | 描述 |
106|---|---|---|112|---|---|---|
107-| dst | 输出 | 目的矢量数据寄存器。dtype必须与`src`一致，搬入VL长度数据。 |113+| dst | 输出 | 目的矢量数据寄存器。仅无返回值类型接口包含该参数。dtype必须与`src`一致，搬入VL长度数据。 |
108| src | 输入 | 源UB地址，实际读取地址必须按dtype对齐。 |114| src | 输入 | 源UB地址，实际读取地址必须按dtype对齐。 |
109 115 
110### 立即数偏移搬入模式116### 立即数偏移搬入模式
@@ -131,7 +137,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem(vector_half& dst,
131 137 
132## 返回值说明138## 返回值说明
133 139 
134-无140+对于返回值类型接口，返回保存广播搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
135 141 
136## 约束说明142## 约束说明
137 143 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_elem2datablock.md+10-4
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中16字节对齐（b16类型）或32字节对齐（b32类型）的起始地址连续读取8个元素，并将每个元素广播到目的矢量数据寄存器对应的一个`DataBlock`（32字节）中。搬运过程中数据格式和内容保持不变。本接口提供三种功能模式：29+从Unified Buffer（UB）中16字节对齐（b16类型）或32字节对齐（b32类型）的起始地址连续读取8个元素，并将每个元素广播到输出结果对应的一个`DataBlock`（32字节）中。结果通过函数返回值返回或写入目的矢量数据寄存器，搬运过程中数据格式和内容保持不变。本接口提供三种功能模式：
30 30 
31- **对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器，由用户自行更新源地址。31- **对齐搬入模式**：将UB源地址的数据搬入到矢量数据寄存器，由用户自行更新源地址。
32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
@@ -39,6 +39,10 @@
39### 对齐搬入模式39### 对齐搬入模式
40 40 
41```c41```c
42+// 通过函数返回值返回结果。
43+__simd_callee__ inline vector_<dtype> asc_loadalign_brc_elem2datablock(__ubuf__ <dtype>* src)
44+ 
45+// 通过引用参数输出结果。
42__simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_<dtype>& dst,46__simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_<dtype>& dst,
43 __ubuf__ <dtype>* src)47 __ubuf__ <dtype>* src)
44```48```
@@ -51,8 +55,10 @@ dtype支持的数据类型为`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`i
51 55 
52```c56```c
53// 示例：half类型。57// 示例：half类型。
58+__simd_callee__ inline vector_half asc_loadalign_brc_elem2datablock(__ubuf__ half* src)
59+ 
54__simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_half& dst,60__simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_half& dst,
55- __ubuf__ half* src)61+ __ubuf__ half* src)
56```62```
57 63 
58### 立即数偏移搬入模式64### 立即数偏移搬入模式
@@ -105,7 +111,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_half& dst,
105 111 
106| 参数名 | 输入/输出 | 描述 |112| 参数名 | 输入/输出 | 描述 |
107|---|---|---|113|---|---|---|
108-| dst | 输出 | 目的矢量数据寄存器。dtype必须与`src`一致，搬入VL长度数据。 |114+| dst | 输出 | 目的矢量数据寄存器。仅无返回值类型接口包含该参数。dtype必须与`src`一致，搬入VL长度数据。 |
109| src | 输入 | 源UB地址，实际读取地址必须满足对齐要求：b16类型必须按16字节对齐，b32类型必须按32字节对齐。 |115| src | 输入 | 源UB地址，实际读取地址必须满足对齐要求：b16类型必须按16字节对齐，b32类型必须按32字节对齐。 |
110 116 
111### 立即数偏移搬入模式117### 立即数偏移搬入模式
@@ -132,7 +138,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_half& dst,
132 138 
133## 返回值说明139## 返回值说明
134 140 
135-无141+对于返回值类型接口，返回保存广播搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
136 142 
137## 约束说明143## 约束说明
138 144 
Adocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_datablock_strided.md+171-0
@@ -0,0 +1,171 @@
1+# asc_loadalign_datablock_strided
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+从Unified Buffer（UB）中32字节对齐的起始地址非连续搬入8个`DataBlock`，并通过函数返回值返回保存搬入结果的矢量数据寄存器。每个`DataBlock`的数据量为32字节，支持配置相邻数据块之间的地址步长和本次搬入的起始读取位置。
30+ 
31+本接口与[asc_loadalign](asc_loadalign.md)的非连续对齐搬入模式功能相同，区别在于本接口通过函数返回值返回结果。
32+ 
33+本接口仅在AIV上生效，非AIV调用直接返回。
34+ 
35+## 函数原型
36+ 
37+```c
38+__simd_callee__ inline vector_<dtype> asc_loadalign_datablock_strided(__ubuf__ <dtype>* src,
39+ uint16_t block_stride,
40+ uint16_t repeat_stride,
41+ vector_bool mask)
42+```
43+ 
44+### dtype支持数据类型
45+ 
46+dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`、`fp4x2_e1m2_t`、`hifloat8_t`、`fp8_e8m0_t`、`fp8_e5m2_t`、`fp8_e4m3fn_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。当dtype为`int4b_t`时，返回值的实际类型为`vector_int4x2_t`。
47+ 
48+### 函数原型典型示例
49+ 
50+```c
51+// 示例：uint8_t类型。
52+__simd_callee__ inline vector_uint8_t asc_loadalign_datablock_strided(__ubuf__ uint8_t* src,
53+ uint16_t block_stride,
54+ uint16_t repeat_stride,
55+ vector_bool mask)
56+```
57+ 
58+## 参数说明
59+ 
60+**表1** 参数说明
61+ 
62+| 参数名 | 输入/输出 | 描述 |
63+|---|---|---|
64+| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |
65+| block_stride | 输入 | 源操作数相邻`DataBlock`之间起始地址的步长，单位为32字节。 |
66+| repeat_stride | 输入 | 本次搬入的起始读取地址相对`src`的偏移，单位为32字节。实际起始读取地址为`src`偏移`repeat_stride × 32`字节。 |
67+| mask | 输入 | 掩码寄存器，用于指示在计算过程中哪些元素参与计算。该接口以`DataBlock`为数据搬运单元。<br>&bull; 当`DataBlock`中的任意一个元素被`mask`筛选成有效元素时，该`DataBlock`中所有数据都会搬入至矢量数据寄存器。<br>&bull; 当`DataBlock`中所有元素都被`mask`筛选成无效元素时，该`DataBlock`中的数据不会搬入到矢量数据寄存器，对应位置的元素设置为0，即使UB越界也不会报错。 |
68+ 
69+矢量数据寄存器和掩码寄存器的详细说明请参见[Reg数据类型定义](../../defs/type/data_type_definition.md)。
70+ 
71+## 返回值说明
72+ 
73+返回保存非连续对齐搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
74+ 
75+## 约束说明
76+ 
77+- 本接口仅在AIV上生效，非AIV调用直接返回。
78+- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用。
79+- 实际读取地址必须按32字节对齐，且有效`DataBlock`的读取范围必须在UB地址空间内且不越界，否则会报错。
80+- 当一个`DataBlock`中的元素全部被`mask`设置为无效时，该`DataBlock`即使越界也不会报错。
81+- `mask`需通过[掩码设置接口](../../defs/type/data_type_definition.md#掩码寄存器)预先赋值后再传入。
82+- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。
83+ 
84+## 调用示例
85+ 
86+将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
87+ 
88+<!-- npu="950" id8 -->
89+ 
90+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下：
91+ 
92+```bash
93+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
94+```
95+<!-- end id8 -->
96+ 
97+```cpp
98+#include <cstdint>
99+#include <iostream>
100+#include <vector>
101+#include "c_api/asc_simd.h"
102+#include "acl/acl.h"
103+ 
104+namespace {
105+constexpr uint32_t DATABLOCK_BYTES = 32;
106+constexpr uint32_t DATABLOCK_COUNT = 8;
107+constexpr uint16_t BLOCK_STRIDE = 2;
108+constexpr uint16_t REPEAT_STRIDE = 1;
109+constexpr uint32_t INPUT_BYTES = 16 * DATABLOCK_BYTES;
110+constexpr uint32_t OUTPUT_BYTES = DATABLOCK_COUNT * DATABLOCK_BYTES;
111+ 
112+__simd_vf__ inline void asc_loadalign_datablock_strided_vf(
113+ __ubuf__ uint8_t* output, __ubuf__ uint8_t* input)
114+{
115+ vector_bool mask = asc_create_mask_b8(PAT_ALL);
116+ vector_uint8_t dst = asc_loadalign_datablock_strided(input, BLOCK_STRIDE, REPEAT_STRIDE, mask);
117+ asc_storealign(output, dst, mask);
118+}
119+ 
120+__global__ __vector__ void asc_loadalign_datablock_strided_kernel(
121+ __gm__ uint8_t* output, __gm__ uint8_t* input)
122+{
123+ asc_init();
124+ __ubuf__ uint8_t input_local[INPUT_BYTES];
125+ __ubuf__ uint8_t output_local[OUTPUT_BYTES];
126+ asc_copy_gm2ub_align(input_local, input, INPUT_BYTES);
127+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
128+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
129+ asc_vf_call<asc_loadalign_datablock_strided_vf>(output_local, input_local);
130+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
131+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
132+ asc_copy_ub2gm_align(output, output_local, OUTPUT_BYTES);
133+ asc_sync();
134+}
135+} // namespace
136+ 
137+int main()
138+{
139+ std::vector<uint8_t> input(INPUT_BYTES);
140+ std::vector<uint8_t> output(OUTPUT_BYTES, 0xff);
141+ std::vector<uint8_t> golden(OUTPUT_BYTES, 0);
142+ for (uint32_t i = 0; i < INPUT_BYTES; ++i) input[i] = static_cast<uint8_t>(i % 251 + 1);
143+ for (uint32_t block = 0; block < DATABLOCK_COUNT; ++block) {
144+ const uint32_t src_offset = (REPEAT_STRIDE + block * BLOCK_STRIDE) * DATABLOCK_BYTES;
145+ const uint32_t dst_offset = block * DATABLOCK_BYTES;
146+ for (uint32_t i = 0; i < DATABLOCK_BYTES; ++i) {
147+ golden[dst_offset + i] = input[src_offset + i];
148+ }
149+ }
150+ 
151+ aclInit(nullptr);
152+ aclrtSetDevice(0);
153+ uint8_t* input_device = nullptr;
154+ uint8_t* output_device = nullptr;
155+ aclrtMalloc(reinterpret_cast<void**>(&input_device), INPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
156+ aclrtMalloc(reinterpret_cast<void**>(&output_device), OUTPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
157+ aclrtMemcpy(input_device, INPUT_BYTES, input.data(), INPUT_BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
158+ asc_loadalign_datablock_strided_kernel<<<1, 0>>>(output_device, input_device);
159+ aclrtSynchronizeDevice();
160+ aclrtMemcpy(output.data(), OUTPUT_BYTES, output_device, OUTPUT_BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
161+ 
162+ const bool passed = output == golden;
163+ std::cout << (passed ? "[Success]" : "[Failed]")
164+ << " asc_loadalign_datablock_strided example." << std::endl;
165+ aclrtFree(input_device);
166+ aclrtFree(output_device);
167+ aclrtResetDevice(0);
168+ aclFinalize();
169+ return passed ? 0 : 1;
170+}
171+```
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_downsample.md+10-3
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中32字节对齐的起始地址读取连续数据并进行2倍下采样。29+从Unified Buffer（UB）中32字节对齐的起始地址读取连续数据并进行2倍下采样，结果通过函数返回值返回或写入目的寄存器。
30 30 
31- 当`dst`为矢量数据寄存器时，读取2×VL（512字节）长度数据，保留偶数下标元素后写入VL长度数据。31- 当`dst`为矢量数据寄存器时，读取2×VL（512字节）长度数据，保留偶数下标元素后写入VL长度数据。
32- 当`dst`为掩码寄存器时，读取VL/4（64字节）长度数据，保留偶数下标bit后写入VL/8长度数据。32- 当`dst`为掩码寄存器时，读取VL/4（64字节）长度数据，保留偶数下标bit后写入VL/8长度数据。
@@ -37,6 +37,8 @@
37- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。37- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
38- **地址寄存器偏移搬入模式**：通过地址寄存器指定相对源起始地址的偏移，常用于Hardware Loop内偏移随循环计数变化的对齐搬入场景。需要与[asc_update_addr_reg](../reg_addr_reg/asc_update_addr_reg.md)配合使用。38- **地址寄存器偏移搬入模式**：通过地址寄存器指定相对源起始地址的偏移，常用于Hardware Loop内偏移随循环计数变化的对齐搬入场景。需要与[asc_update_addr_reg](../reg_addr_reg/asc_update_addr_reg.md)配合使用。
39 39 
40+对齐搬入模式中，通过函数返回值返回掩码寄存器时，请使用[asc_loadalign_mask_downsample](asc_loadalign_mask_downsample.md)。
41+ 
40本接口仅在AIV上生效，非AIV调用直接返回。42本接口仅在AIV上生效，非AIV调用直接返回。
41 43 
42## 函数原型44## 函数原型
@@ -44,6 +46,9 @@
44### 对齐搬入模式46### 对齐搬入模式
45 47 
46```c48```c
49+// 通过函数返回值返回矢量数据寄存器。
50+__simd_callee__ inline vector_<dtype> asc_loadalign_downsample(__ubuf__ <dtype>* src)
51+ 
47// 搬入矢量数据寄存器。52// 搬入矢量数据寄存器。
48__simd_callee__ inline void asc_loadalign_downsample(vector_<dtype>& dst,53__simd_callee__ inline void asc_loadalign_downsample(vector_<dtype>& dst,
49 __ubuf__ <dtype>* src)54 __ubuf__ <dtype>* src)
@@ -60,6 +65,8 @@ dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`
60 65 
61```c66```c
62// 示例：half类型。67// 示例：half类型。
68+__simd_callee__ inline vector_half asc_loadalign_downsample(__ubuf__ half* src)
69+ 
63__simd_callee__ inline void asc_loadalign_downsample(vector_half& dst,70__simd_callee__ inline void asc_loadalign_downsample(vector_half& dst,
64 __ubuf__ half* src)71 __ubuf__ half* src)
65```72```
@@ -124,7 +131,7 @@ __simd_callee__ inline void asc_loadalign_downsample(vector_half& dst,
124 131 
125| 参数名 | 输入/输出 | 描述 |132| 参数名 | 输入/输出 | 描述 |
126|---|---|---|133|---|---|---|
127-| dst | 输出 | 目的矢量数据寄存器或掩码寄存器。<br>&bull; 当`dst`为矢量数据寄存器，dtype必须与`src`一致，搬入VL长度数据。<br>&bull; 当`dst`为掩码寄存器，搬入VL/8长度数据。 |134+| dst | 输出 | 目的矢量数据寄存器或掩码寄存器，仅无返回值类型接口包含该参数。<br>&bull; 当`dst`为矢量数据寄存器，dtype必须与`src`一致，搬入VL长度数据。<br>&bull; 当`dst`为掩码寄存器，搬入VL/8长度数据。 |
128| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |135| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |
129 136 
130### 立即数偏移搬入模式137### 立即数偏移搬入模式
@@ -151,7 +158,7 @@ __simd_callee__ inline void asc_loadalign_downsample(vector_half& dst,
151 158 
152## 返回值说明159## 返回值说明
153 160 
154-无161+对于返回值类型接口，返回保存下采样搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
155 162 
156## 约束说明163## 约束说明
157 164 
Adocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_mask.md+133-0
@@ -0,0 +1,133 @@
1+# asc_loadalign_mask
L
LLycheeeee1 天前

新增接口更新README、c_api.md等索引文件。

likedislike
L
lihuaichao
1 天前 评论：
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+从Unified Buffer（UB）中32字节对齐的起始地址读取VL/8长度数据，并通过函数返回值返回掩码寄存器。搬运过程中数据格式和内容保持不变。连续搬入时，需要在每次调用前手动更新源地址。
30+ 
31+本接口与[asc_loadalign](asc_loadalign.md)连续对齐搬入模式中目的操作数为掩码寄存器的原型功能相同，区别在于本接口通过函数返回值返回结果。
32+ 
33+本接口仅在AIV上生效，非AIV调用直接返回。
34+ 
35+## 函数原型
36+ 
37+```c
38+__simd_callee__ inline vector_bool asc_loadalign_mask(__ubuf__ uint32_t* src)
39+```
40+ 
41+## 参数说明
42+ 
43+**表1** 参数说明
44+ 
45+| 参数名 | 输入/输出 | 描述 |
46+|---|---|---|
47+| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐，搬入VL/8长度数据。 |
48+ 
49+掩码寄存器的详细说明请参见[Reg数据类型定义](../../defs/type/data_type_definition.md)。
50+ 
51+## 返回值说明
52+ 
53+返回保存连续对齐搬入结果的掩码寄存器，类型为`vector_bool`。
54+ 
55+## 约束说明
56+ 
57+- 本接口仅在AIV上生效，非AIV调用直接返回。
58+- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用。
59+- `src`的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
60+- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。
61+ 
62+## 调用示例
63+ 
64+将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65+ 
66+<!-- npu="950" id8 -->
67+ 
68+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下：
69+ 
70+```bash
71+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
72+```
73+<!-- end id8 -->
74+ 
75+```cpp
76+#include <cstdint>
77+#include <iostream>
78+#include <vector>
79+#include "c_api/asc_simd.h"
80+#include "acl/acl.h"
81+ 
82+namespace {
83+constexpr uint32_t MASK_BYTES = 32;
84+ 
85+__simd_vf__ inline void asc_loadalign_mask_vf(
86+ __ubuf__ uint32_t* output, __ubuf__ uint32_t* mask_input)
87+{
88+ vector_bool mask = asc_loadalign_mask(mask_input);
89+ asc_storealign(output, mask);
90+}
91+ 
92+__global__ __vector__ void asc_loadalign_mask_kernel(
93+ __gm__ uint32_t* output, __gm__ uint32_t* mask_input)
94+{
95+ asc_init();
96+ __ubuf__ uint32_t output_local[MASK_BYTES / sizeof(uint32_t)];
97+ __ubuf__ uint32_t mask_local[MASK_BYTES / sizeof(uint32_t)];
98+ asc_copy_gm2ub_align(mask_local, mask_input, MASK_BYTES);
99+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
100+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
101+ asc_vf_call<asc_loadalign_mask_vf>(output_local, mask_local);
102+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
103+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
104+ asc_copy_ub2gm_align(output, output_local, MASK_BYTES);
105+ asc_sync();
106+}
107+} // namespace
108+ 
109+int main()
110+{
111+ std::vector<uint32_t> output(MASK_BYTES / sizeof(uint32_t), 0);
112+ std::vector<uint32_t> mask_input(MASK_BYTES / sizeof(uint32_t), 0x55555555u);
113+ 
114+ aclInit(nullptr);
115+ aclrtSetDevice(0);
116+ uint32_t* output_device = nullptr;
117+ uint32_t* mask_device = nullptr;
118+ aclrtMalloc(reinterpret_cast<void**>(&output_device), MASK_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
119+ aclrtMalloc(reinterpret_cast<void**>(&mask_device), MASK_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
120+ aclrtMemcpy(mask_device, MASK_BYTES, mask_input.data(), MASK_BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
121+ asc_loadalign_mask_kernel<<<1, 0>>>(output_device, mask_device);
122+ aclrtSynchronizeDevice();
123+ aclrtMemcpy(output.data(), MASK_BYTES, output_device, MASK_BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
124+ 
125+ const bool passed = output == mask_input;
126+ std::cout << (passed ? "[Success]" : "[Failed]") << " asc_loadalign_mask example." << std::endl;
127+ aclrtFree(output_device);
128+ aclrtFree(mask_device);
129+ aclrtResetDevice(0);
130+ aclFinalize();
131+ return passed ? 0 : 1;
132+}
133+```
Adocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_mask_downsample.md+136-0
@@ -0,0 +1,136 @@
1+# asc_loadalign_mask_downsample
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+从Unified Buffer（UB）中32字节对齐的起始地址读取VL/4长度数据，保留偶数下标bit，得到VL/8长度数据，并通过函数返回值返回掩码寄存器。
30+ 
31+本接口与[asc_loadalign_downsample](asc_loadalign_downsample.md)对齐搬入模式中目的操作数为掩码寄存器的原型功能相同，区别在于本接口通过函数返回值返回结果。
32+ 
33+本接口仅在AIV上生效，非AIV调用直接返回。
34+ 
35+## 函数原型
36+ 
37+```c
38+__simd_callee__ inline vector_bool asc_loadalign_mask_downsample(__ubuf__ uint32_t* src)
39+```
40+ 
41+## 参数说明
42+ 
43+**表1** 参数说明
44+ 
45+| 参数名 | 输入/输出 | 描述 |
46+|---|---|---|
47+| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐，读取VL/4长度数据。 |
48+ 
49+掩码寄存器的详细说明请参见[Reg数据类型定义](../../defs/type/data_type_definition.md)。
M
Mmunanhw1 天前

参数里没有掩码寄存器

likedislike
50+ 
51+## 返回值说明
52+ 
53+返回保存2倍下采样搬入结果的掩码寄存器，类型为`vector_bool`，有效数据长度为VL/8。
54+ 
55+## 约束说明
56+ 
57+- 本接口仅在AIV上生效，非AIV调用直接返回。
58+- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用。
59+- `src`的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
60+- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。
61+ 
62+## 调用示例
63+ 
64+将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65+ 
66+<!-- npu="950" id8 -->
67+ 
68+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下：
69+ 
70+```bash
71+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
72+```
73+<!-- end id8 -->
74+ 
75+```cpp
76+#include <cstdint>
77+#include <iostream>
78+#include <vector>
79+#include "c_api/asc_simd.h"
80+#include "acl/acl.h"
81+ 
82+namespace {
83+constexpr uint32_t MASK_INPUT_BYTES = 64;
84+constexpr uint32_t MASK_OUTPUT_BYTES = 32;
85+ 
86+__simd_vf__ inline void asc_loadalign_mask_downsample_vf(
87+ __ubuf__ uint32_t* output, __ubuf__ uint32_t* mask_input)
88+{
89+ vector_bool mask = asc_loadalign_mask_downsample(mask_input);
90+ asc_storealign(output, mask);
91+}
92+ 
93+__global__ __vector__ void asc_loadalign_mask_downsample_kernel(
94+ __gm__ uint32_t* output, __gm__ uint32_t* mask_input)
95+{
96+ asc_init();
97+ __ubuf__ uint32_t output_local[MASK_OUTPUT_BYTES / sizeof(uint32_t)];
98+ __ubuf__ uint32_t mask_local[MASK_INPUT_BYTES / sizeof(uint32_t)];
99+ asc_copy_gm2ub_align(mask_local, mask_input, MASK_INPUT_BYTES);
100+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
101+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
102+ asc_vf_call<asc_loadalign_mask_downsample_vf>(output_local, mask_local);
103+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
104+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
105+ asc_copy_ub2gm_align(output, output_local, MASK_OUTPUT_BYTES);
106+ asc_sync();
107+}
108+} // namespace
109+ 
110+int main()
111+{
112+ std::vector<uint32_t> output(MASK_OUTPUT_BYTES / sizeof(uint32_t), 0);
113+ std::vector<uint32_t> golden(MASK_OUTPUT_BYTES / sizeof(uint32_t), 0x55555555u);
114+ std::vector<uint32_t> mask_input(MASK_INPUT_BYTES / sizeof(uint32_t), 0x33333333u);
115+ 
116+ aclInit(nullptr);
117+ aclrtSetDevice(0);
118+ uint32_t* output_device = nullptr;
119+ uint32_t* mask_device = nullptr;
120+ aclrtMalloc(reinterpret_cast<void**>(&output_device), MASK_OUTPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
121+ aclrtMalloc(reinterpret_cast<void**>(&mask_device), MASK_INPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
122+ aclrtMemcpy(mask_device, MASK_INPUT_BYTES, mask_input.data(), MASK_INPUT_BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
123+ asc_loadalign_mask_downsample_kernel<<<1, 0>>>(output_device, mask_device);
124+ aclrtSynchronizeDevice();
125+ aclrtMemcpy(output.data(), MASK_OUTPUT_BYTES, output_device, MASK_OUTPUT_BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
126+ 
127+ const bool passed = output == golden;
128+ std::cout << (passed ? "[Success]" : "[Failed]")
129+ << " asc_loadalign_mask_downsample example." << std::endl;
130+ aclrtFree(output_device);
131+ aclrtFree(mask_device);
132+ aclrtResetDevice(0);
133+ aclFinalize();
134+ return passed ? 0 : 1;
135+}
136+```
Adocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_mask_upsample.md+140-0
@@ -0,0 +1,140 @@
1+# asc_loadalign_mask_upsample
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+从Unified Buffer（UB）中16字节对齐的起始地址读取VL/16长度数据，将每个bit重复两次，得到VL/8长度数据，并通过函数返回值返回掩码寄存器。
30+ 
31+本接口与[asc_loadalign_upsample](asc_loadalign_upsample.md)对齐搬入模式中目的操作数为掩码寄存器的原型功能相同，区别在于本接口通过函数返回值返回结果。
32+ 
33+本接口仅在AIV上生效，非AIV调用直接返回。
34+ 
35+## 函数原型
36+ 
37+```c
38+__simd_callee__ inline vector_bool asc_loadalign_mask_upsample(__ubuf__ uint32_t* src)
39+```
40+ 
41+## 参数说明
42+ 
43+**表1** 参数说明
44+ 
45+| 参数名 | 输入/输出 | 描述 |
46+|---|---|---|
47+| src | 输入 | 源UB地址，实际读取地址必须按16字节对齐，读取VL/16长度数据。 |
48+ 
49+掩码寄存器的详细说明请参见[Reg数据类型定义](../../defs/type/data_type_definition.md)。
M
Mmunanhw1 天前

参数里没有掩码寄存器

likedislike
50+ 
51+## 返回值说明
52+ 
53+返回保存2倍上采样搬入结果的掩码寄存器，类型为`vector_bool`，有效数据长度为VL/8。
54+ 
55+## 约束说明
56+ 
57+- 本接口仅在AIV上生效，非AIV调用直接返回。
58+- 本接口在Vector Function（`__simd_vf__`标记的函数）内调用。
59+- `src`的实际读取地址必须按16字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
60+- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。
61+ 
62+## 调用示例
63+ 
64+将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65+ 
66+<!-- npu="950" id8 -->
67+ 
68+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下：
69+ 
70+```bash
71+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
72+```
73+<!-- end id8 -->
74+ 
75+```cpp
76+#include <cstdint>
77+#include <iostream>
78+#include <vector>
79+#include "c_api/asc_simd.h"
80+#include "acl/acl.h"
81+ 
82+namespace {
83+constexpr uint32_t MASK_DATA_BYTES = 16;
84+constexpr uint32_t MASK_STORAGE_BYTES = 32;
85+constexpr uint32_t MASK_OUTPUT_BYTES = 32;
86+ 
87+__simd_vf__ inline void asc_loadalign_mask_upsample_vf(
88+ __ubuf__ uint32_t* output, __ubuf__ uint32_t* mask_input)
89+{
90+ vector_bool mask = asc_loadalign_mask_upsample(mask_input);
91+ asc_storealign(output, mask);
92+}
93+ 
94+__global__ __vector__ void asc_loadalign_mask_upsample_kernel(
95+ __gm__ uint32_t* output, __gm__ uint32_t* mask_input)
96+{
97+ asc_init();
98+ __ubuf__ uint32_t output_local[MASK_OUTPUT_BYTES / sizeof(uint32_t)];
99+ __ubuf__ uint32_t mask_local[MASK_STORAGE_BYTES / sizeof(uint32_t)];
100+ asc_copy_gm2ub_align(mask_local, mask_input, MASK_DATA_BYTES);
101+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
102+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
103+ asc_vf_call<asc_loadalign_mask_upsample_vf>(output_local, mask_local);
104+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
105+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
106+ asc_copy_ub2gm_align(output, output_local, MASK_OUTPUT_BYTES);
107+ asc_sync();
108+}
109+} // namespace
110+ 
111+int main()
112+{
113+ std::vector<uint32_t> output(MASK_OUTPUT_BYTES / sizeof(uint32_t), 0);
114+ std::vector<uint32_t> golden(MASK_OUTPUT_BYTES / sizeof(uint32_t), 0x33333333u);
115+ std::vector<uint32_t> mask_input(MASK_STORAGE_BYTES / sizeof(uint32_t), 0);
116+ for (uint32_t i = 0; i < MASK_DATA_BYTES / sizeof(uint32_t); ++i) {
117+ mask_input[i] = 0x55555555u;
118+ }
119+ 
120+ aclInit(nullptr);
121+ aclrtSetDevice(0);
122+ uint32_t* output_device = nullptr;
123+ uint32_t* mask_device = nullptr;
124+ aclrtMalloc(reinterpret_cast<void**>(&output_device), MASK_OUTPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
125+ aclrtMalloc(reinterpret_cast<void**>(&mask_device), MASK_STORAGE_BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
126+ aclrtMemcpy(mask_device, MASK_STORAGE_BYTES, mask_input.data(), MASK_STORAGE_BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
127+ asc_loadalign_mask_upsample_kernel<<<1, 0>>>(output_device, mask_device);
128+ aclrtSynchronizeDevice();
129+ aclrtMemcpy(output.data(), MASK_OUTPUT_BYTES, output_device, MASK_OUTPUT_BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
130+ 
131+ const bool passed = output == golden;
132+ std::cout << (passed ? "[Success]" : "[Failed]")
133+ << " asc_loadalign_mask_upsample example." << std::endl;
134+ aclrtFree(output_device);
135+ aclrtFree(mask_device);
136+ aclrtResetDevice(0);
137+ aclFinalize();
138+ return passed ? 0 : 1;
139+}
140+```
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_unpack.md+9-3
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中32字节对齐的起始地址读取VL/2长度的连续数据，在每个源元素后补1个值为0的元素，并将得到的VL长度数据写入目的矢量数据寄存器。本接口提供三种功能模式：29+从Unified Buffer（UB）中32字节对齐的起始地址读取VL/2长度的连续数据，在每个源元素后补1个值为0的元素，并将得到的VL长度数据通过函数返回值返回或写入目的矢量数据寄存器。本接口提供三种功能模式：
30 30 
31- **对齐搬入模式**：将UB源地址的数据解包后搬入到矢量数据寄存器，由用户自行更新源地址。31- **对齐搬入模式**：将UB源地址的数据解包后搬入到矢量数据寄存器，由用户自行更新源地址。
32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
@@ -39,6 +39,10 @@
39### 对齐搬入模式39### 对齐搬入模式
40 40 
41```c41```c
42+// 通过函数返回值返回结果。
43+__simd_callee__ inline vector_<dtype> asc_loadalign_unpack(__ubuf__ <dtype>* src)
44+ 
45+// 通过引用参数输出结果。
42__simd_callee__ inline void asc_loadalign_unpack(vector_<dtype>& dst,46__simd_callee__ inline void asc_loadalign_unpack(vector_<dtype>& dst,
43 __ubuf__ <dtype>* src)47 __ubuf__ <dtype>* src)
44```48```
@@ -51,6 +55,8 @@ dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`
51 55 
52```c56```c
53// 示例：half类型。57// 示例：half类型。
58+__simd_callee__ inline vector_half asc_loadalign_unpack(__ubuf__ half* src)
59+ 
54__simd_callee__ inline void asc_loadalign_unpack(vector_half& dst,60__simd_callee__ inline void asc_loadalign_unpack(vector_half& dst,
55 __ubuf__ half* src)61 __ubuf__ half* src)
56```62```
@@ -105,7 +111,7 @@ __simd_callee__ inline void asc_loadalign_unpack(vector_half& dst,
105 111 
106| 参数名 | 输入/输出 | 描述 |112| 参数名 | 输入/输出 | 描述 |
107|---|---|---|113|---|---|---|
108-| dst | 输出 | 目的矢量数据寄存器。dtype必须与`src`一致，每个源元素后补1个值为0的元素，搬入VL长度数据。 |114+| dst | 输出 | 目的矢量数据寄存器。仅无返回值类型接口包含该参数。dtype必须与`src`一致，每个源元素后补1个值为0的元素，搬入VL长度数据。 |
109| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |115| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |
110 116 
111### 立即数偏移搬入模式117### 立即数偏移搬入模式
@@ -132,7 +138,7 @@ __simd_callee__ inline void asc_loadalign_unpack(vector_half& dst,
132 138 
133## 返回值说明139## 返回值说明
134 140 
135-无141+对于返回值类型接口，返回保存解包搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
136 142 
137## 约束说明143## 约束说明
138 144 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_unpack4.md+9-3
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中32字节对齐的起始地址读取VL/4长度的连续数据，在每个源元素后补3个值为0的元素，并将得到的VL长度数据写入目的矢量数据寄存器。本接口提供三种功能模式：29+从Unified Buffer（UB）中32字节对齐的起始地址读取VL/4长度的连续数据，在每个源元素后补3个值为0的元素，并将得到的VL长度数据通过函数返回值返回或写入目的矢量数据寄存器。本接口提供三种功能模式：
30 30 
31- **对齐搬入模式**：将UB源地址的数据解包后搬入到矢量数据寄存器，由用户自行更新源地址。31- **对齐搬入模式**：将UB源地址的数据解包后搬入到矢量数据寄存器，由用户自行更新源地址。
32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。32- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
@@ -39,6 +39,10 @@
39### 对齐搬入模式39### 对齐搬入模式
40 40 
41```c41```c
42+// 通过函数返回值返回结果。
43+__simd_callee__ inline vector_<dtype> asc_loadalign_unpack4(__ubuf__ <dtype>* src)
44+ 
45+// 通过引用参数输出结果。
42__simd_callee__ inline void asc_loadalign_unpack4(vector_<dtype>& dst,46__simd_callee__ inline void asc_loadalign_unpack4(vector_<dtype>& dst,
43 __ubuf__ <dtype>* src)47 __ubuf__ <dtype>* src)
44```48```
@@ -51,6 +55,8 @@ dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`
51 55 
52```c56```c
53// 示例：int8_t类型。57// 示例：int8_t类型。
58+__simd_callee__ inline vector_int8_t asc_loadalign_unpack4(__ubuf__ int8_t* src)
59+ 
54__simd_callee__ inline void asc_loadalign_unpack4(vector_int8_t& dst,60__simd_callee__ inline void asc_loadalign_unpack4(vector_int8_t& dst,
55 __ubuf__ int8_t* src)61 __ubuf__ int8_t* src)
56```62```
@@ -105,7 +111,7 @@ __simd_callee__ inline void asc_loadalign_unpack4(vector_int8_t& dst,
105 111 
106| 参数名 | 输入/输出 | 描述 |112| 参数名 | 输入/输出 | 描述 |
107|---|---|---|113|---|---|---|
108-| dst | 输出 | 目的矢量数据寄存器。dtype必须与`src`一致，每个源元素后补3个值为0的元素，搬入VL长度数据。 |114+| dst | 输出 | 目的矢量数据寄存器。仅无返回值类型接口包含该参数。dtype必须与`src`一致，每个源元素后补3个值为0的元素，搬入VL长度数据。 |
109| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |115| src | 输入 | 源UB地址，实际读取地址必须按32字节对齐。 |
110 116 
111### 立即数偏移搬入模式117### 立即数偏移搬入模式
@@ -132,7 +138,7 @@ __simd_callee__ inline void asc_loadalign_unpack4(vector_int8_t& dst,
132 138 
133## 返回值说明139## 返回值说明
134 140 
135-无141+对于返回值类型接口，返回保存解包搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
136 142 
137## 约束说明143## 约束说明
138 144 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_upsample.md+10-3
@@ -26,7 +26,7 @@
26 26 
27## 功能说明27## 功能说明
28 28 
29-从Unified Buffer（UB）中满足对应对齐要求的起始地址读取连续数据并进行2倍上采样。29+从Unified Buffer（UB）中满足对应对齐要求的起始地址读取连续数据并进行2倍上采样，结果通过函数返回值返回或写入目的寄存器。
30 30 
31- 当`dst`为矢量数据寄存器时，从UB中32字节对齐的起始地址读取VL/2长度数据，将每个源元素重复两次后写入VL长度数据。31- 当`dst`为矢量数据寄存器时，从UB中32字节对齐的起始地址读取VL/2长度数据，将每个源元素重复两次后写入VL长度数据。
32- 当`dst`为掩码寄存器时，从UB中16字节对齐的起始地址读取VL/16长度数据，将每个bit重复两次后写入VL/8长度数据。32- 当`dst`为掩码寄存器时，从UB中16字节对齐的起始地址读取VL/16长度数据，将每个bit重复两次后写入VL/8长度数据。
@@ -37,6 +37,8 @@
37- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。37- **立即数偏移搬入模式**：从相对源起始地址偏移指定距离的位置搬入数据。本接口不会自动更新源地址。
38- **地址寄存器偏移搬入模式**：通过地址寄存器指定相对源起始地址的偏移，常用于Hardware Loop内偏移随循环计数变化的对齐搬入场景。需要与[asc_update_addr_reg](../reg_addr_reg/asc_update_addr_reg.md)配合使用。38- **地址寄存器偏移搬入模式**：通过地址寄存器指定相对源起始地址的偏移，常用于Hardware Loop内偏移随循环计数变化的对齐搬入场景。需要与[asc_update_addr_reg](../reg_addr_reg/asc_update_addr_reg.md)配合使用。
39 39 
40+对齐搬入模式中，通过函数返回值返回掩码寄存器时，请使用[asc_loadalign_mask_upsample](asc_loadalign_mask_upsample.md)。
41+ 
40本接口仅在AIV上生效，非AIV调用直接返回。42本接口仅在AIV上生效，非AIV调用直接返回。
41 43 
42## 函数原型44## 函数原型
@@ -44,6 +46,9 @@
44### 对齐搬入模式46### 对齐搬入模式
45 47 
46```c48```c
49+// 通过函数返回值返回矢量数据寄存器。
50+__simd_callee__ inline vector_<dtype> asc_loadalign_upsample(__ubuf__ <dtype>* src)
51+ 
47// 搬入矢量数据寄存器。52// 搬入矢量数据寄存器。
48__simd_callee__ inline void asc_loadalign_upsample(vector_<dtype>& dst,53__simd_callee__ inline void asc_loadalign_upsample(vector_<dtype>& dst,
49 __ubuf__ <dtype>* src)54 __ubuf__ <dtype>* src)
@@ -60,6 +65,8 @@ dtype支持的数据类型为`int4b_t`、`int8_t`、`uint8_t`、`fp4x2_e2m1_t`
60 65 
61```c66```c
62// 示例：half类型。67// 示例：half类型。
68+__simd_callee__ inline vector_half asc_loadalign_upsample(__ubuf__ half* src)
69+ 
63__simd_callee__ inline void asc_loadalign_upsample(vector_half& dst,70__simd_callee__ inline void asc_loadalign_upsample(vector_half& dst,
64 __ubuf__ half* src)71 __ubuf__ half* src)
65```72```
@@ -124,7 +131,7 @@ __simd_callee__ inline void asc_loadalign_upsample(vector_half& dst,
124 131 
125| 参数名 | 输入/输出 | 描述 |132| 参数名 | 输入/输出 | 描述 |
126|---|---|---|133|---|---|---|
127-| dst | 输出 | 目的矢量数据寄存器或掩码寄存器。<br>&bull; 当`dst`为矢量数据寄存器，dtype必须与`src`一致，搬入VL长度数据。<br>&bull; 当`dst`为掩码寄存器，搬入VL/8长度数据。 |134+| dst | 输出 | 目的矢量数据寄存器或掩码寄存器，仅无返回值类型接口包含该参数。<br>&bull; 当`dst`为矢量数据寄存器，dtype必须与`src`一致，搬入VL长度数据。<br>&bull; 当`dst`为掩码寄存器，搬入VL/8长度数据。 |
128| src | 输入 | 源UB地址。<br>&bull; 当`dst`为矢量数据寄存器，实际读取地址必须按32字节对齐。<br>&bull; 当`dst`为掩码寄存器，实际读取地址必须按16字节对齐。 |135| src | 输入 | 源UB地址。<br>&bull; 当`dst`为矢量数据寄存器，实际读取地址必须按32字节对齐。<br>&bull; 当`dst`为掩码寄存器，实际读取地址必须按16字节对齐。 |
129 136 
130### 立即数偏移搬入模式137### 立即数偏移搬入模式
@@ -151,7 +158,7 @@ __simd_callee__ inline void asc_loadalign_upsample(vector_half& dst,
151 158 
152## 返回值说明159## 返回值说明
153 160 
154-无161+对于返回值类型接口，返回保存上采样搬入结果的矢量数据寄存器，数据类型与`src`保持一致。
155 162 
156## 约束说明163## 约束说明
157 164 
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/load.md+4-0
@@ -3,6 +3,8 @@
3## Reg对齐搬入3## Reg对齐搬入
4 4 
5- **[asc_loadalign](asc_loadalign.md)**5- **[asc_loadalign](asc_loadalign.md)**
6+- **[asc_loadalign_datablock_strided](asc_loadalign_datablock_strided.md)**
7+- **[asc_loadalign_mask](asc_loadalign_mask.md)**
6- **[asc_loadalign_brc_datablock](asc_loadalign_brc_datablock.md)**8- **[asc_loadalign_brc_datablock](asc_loadalign_brc_datablock.md)**
7- **[asc_loadalign_brc_datablock_postupdate](asc_loadalign_brc_datablock_postupdate.md)**9- **[asc_loadalign_brc_datablock_postupdate](asc_loadalign_brc_datablock_postupdate.md)**
8- **[asc_loadalign_brc_elem](asc_loadalign_brc_elem.md)**10- **[asc_loadalign_brc_elem](asc_loadalign_brc_elem.md)**
@@ -12,6 +14,7 @@
12- **[asc_loadalign_deintlv](asc_loadalign_deintlv.md)**14- **[asc_loadalign_deintlv](asc_loadalign_deintlv.md)**
13- **[asc_loadalign_deintlv_postupdate](asc_loadalign_deintlv_postupdate.md)**15- **[asc_loadalign_deintlv_postupdate](asc_loadalign_deintlv_postupdate.md)**
14- **[asc_loadalign_downsample](asc_loadalign_downsample.md)**16- **[asc_loadalign_downsample](asc_loadalign_downsample.md)**
17+- **[asc_loadalign_mask_downsample](asc_loadalign_mask_downsample.md)**
15- **[asc_loadalign_downsample_postupdate](asc_loadalign_downsample_postupdate.md)**18- **[asc_loadalign_downsample_postupdate](asc_loadalign_downsample_postupdate.md)**
16- **[asc_loadalign_postupdate](asc_loadalign_postupdate.md)**19- **[asc_loadalign_postupdate](asc_loadalign_postupdate.md)**
17- **[asc_loadalign_unpack](asc_loadalign_unpack.md)**20- **[asc_loadalign_unpack](asc_loadalign_unpack.md)**
@@ -19,6 +22,7 @@
19- **[asc_loadalign_unpack4_postupdate](asc_loadalign_unpack4_postupdate.md)**22- **[asc_loadalign_unpack4_postupdate](asc_loadalign_unpack4_postupdate.md)**
20- **[asc_loadalign_unpack_postupdate](asc_loadalign_unpack_postupdate.md)**23- **[asc_loadalign_unpack_postupdate](asc_loadalign_unpack_postupdate.md)**
21- **[asc_loadalign_upsample](asc_loadalign_upsample.md)**24- **[asc_loadalign_upsample](asc_loadalign_upsample.md)**
25+- **[asc_loadalign_mask_upsample](asc_loadalign_mask_upsample.md)**
22- **[asc_loadalign_upsample_postupdate](asc_loadalign_upsample_postupdate.md)**26- **[asc_loadalign_upsample_postupdate](asc_loadalign_upsample_postupdate.md)**
23 27 
24## Reg非对齐搬入28## Reg非对齐搬入
Mimpl/c_api/instr_impl/npu_arch_3510/vector_datamove_impl.h+1016-0
@@ -9039,6 +9039,1022 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock_postupdate(
9039 asc_loadalign_brc_elem2datablock_postupdate_impl(dst, src, offset);9039 asc_loadalign_brc_elem2datablock_postupdate_impl(dst, src, offset);
9040}9040}
9041 9041 
9042+// ==========return-value load APIs==========
9043+// ========== return vector data register APIs ==========
9044+// ========== asc_load ==========
9045+__simd_callee__ inline vector_int4x2_t asc_load(__ubuf__ int4b_t* src)
9046+{
9047+ vector_int4x2_t dst;
9048+ asc_load_impl(dst, src);
9049+ return dst;
9050+}
9051+ 
9052+__simd_callee__ inline vector_int8_t asc_load(__ubuf__ int8_t* src)
9053+{
9054+ vector_int8_t dst;
9055+ asc_load_impl(dst, src);
9056+ return dst;
9057+}
9058+ 
9059+__simd_callee__ inline vector_uint8_t asc_load(__ubuf__ uint8_t* src)
9060+{
9061+ vector_uint8_t dst;
9062+ asc_load_impl(dst, src);
9063+ return dst;
9064+}
9065+ 
9066+__simd_callee__ inline vector_fp4x2_e2m1_t asc_load(__ubuf__ fp4x2_e2m1_t* src)
9067+{
9068+ vector_fp4x2_e2m1_t dst;
9069+ asc_load_impl(dst, src);
9070+ return dst;
9071+}
9072+ 
9073+__simd_callee__ inline vector_fp4x2_e1m2_t asc_load(__ubuf__ fp4x2_e1m2_t* src)
9074+{
9075+ vector_fp4x2_e1m2_t dst;
9076+ asc_load_impl(dst, src);
9077+ return dst;
9078+}
9079+ 
9080+__simd_callee__ inline vector_hifloat8_t asc_load(__ubuf__ hifloat8_t* src)
9081+{
9082+ vector_hifloat8_t dst;
9083+ asc_load_impl(dst, src);
9084+ return dst;
9085+}
9086+ 
9087+__simd_callee__ inline vector_fp8_e8m0_t asc_load(__ubuf__ fp8_e8m0_t* src)
9088+{
9089+ vector_fp8_e8m0_t dst;
9090+ asc_load_impl(dst, src);
9091+ return dst;
9092+}
9093+ 
9094+__simd_callee__ inline vector_fp8_e5m2_t asc_load(__ubuf__ fp8_e5m2_t* src)
9095+{
9096+ vector_fp8_e5m2_t dst;
9097+ asc_load_impl(dst, src);
9098+ return dst;
9099+}
9100+ 
9101+__simd_callee__ inline vector_fp8_e4m3fn_t asc_load(__ubuf__ fp8_e4m3fn_t* src)
9102+{
9103+ vector_fp8_e4m3fn_t dst;
9104+ asc_load_impl(dst, src);
9105+ return dst;
9106+}
9107+ 
9108+__simd_callee__ inline vector_int16_t asc_load(__ubuf__ int16_t* src)
9109+{
9110+ vector_int16_t dst;
9111+ asc_load_impl(dst, src);
9112+ return dst;
9113+}
9114+ 
9115+__simd_callee__ inline vector_uint16_t asc_load(__ubuf__ uint16_t* src)
9116+{
9117+ vector_uint16_t dst;
9118+ asc_load_impl(dst, src);
9119+ return dst;
9120+}
9121+ 
9122+__simd_callee__ inline vector_half asc_load(__ubuf__ half* src)
9123+{
9124+ vector_half dst;
9125+ asc_load_impl(dst, src);
9126+ return dst;
9127+}
9128+ 
9129+__simd_callee__ inline vector_bfloat16_t asc_load(__ubuf__ bfloat16_t* src)
9130+{
9131+ vector_bfloat16_t dst;
9132+ asc_load_impl(dst, src);
9133+ return dst;
9134+}
9135+ 
9136+__simd_callee__ inline vector_int32_t asc_load(__ubuf__ int32_t* src)
9137+{
9138+ vector_int32_t dst;
9139+ asc_load_impl(dst, src);
9140+ return dst;
9141+}
9142+ 
9143+__simd_callee__ inline vector_uint32_t asc_load(__ubuf__ uint32_t* src)
9144+{
9145+ vector_uint32_t dst;
9146+ asc_load_impl(dst, src);
9147+ return dst;
9148+}
9149+ 
9150+__simd_callee__ inline vector_float asc_load(__ubuf__ float* src)
9151+{
9152+ vector_float dst;
9153+ asc_load_impl(dst, src);
9154+ return dst;
9155+}
9156+ 
9157+// ========== asc_loadalign ==========
9158+__simd_callee__ inline vector_int4x2_t asc_loadalign(__ubuf__ int4b_t* src)
9159+{
9160+ vector_int4x2_t dst;
9161+ asc_loadalign_impl(dst, src);
9162+ return dst;
9163+}
9164+ 
9165+__simd_callee__ inline vector_int8_t asc_loadalign(__ubuf__ int8_t* src)
9166+{
9167+ vector_int8_t dst;
9168+ asc_loadalign_impl(dst, src);
9169+ return dst;
9170+}
9171+ 
9172+__simd_callee__ inline vector_uint8_t asc_loadalign(__ubuf__ uint8_t* src)
9173+{
9174+ vector_uint8_t dst;
9175+ asc_loadalign_impl(dst, src);
9176+ return dst;
9177+}
9178+ 
9179+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign(__ubuf__ fp4x2_e2m1_t* src)
9180+{
9181+ vector_fp4x2_e2m1_t dst;
9182+ asc_loadalign_impl(dst, src);
9183+ return dst;
9184+}
9185+ 
9186+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign(__ubuf__ fp4x2_e1m2_t* src)
9187+{
9188+ vector_fp4x2_e1m2_t dst;
9189+ asc_loadalign_impl(dst, src);
9190+ return dst;
9191+}
9192+ 
9193+__simd_callee__ inline vector_hifloat8_t asc_loadalign(__ubuf__ hifloat8_t* src)
9194+{
9195+ vector_hifloat8_t dst;
9196+ asc_loadalign_impl(dst, src);
9197+ return dst;
9198+}
9199+ 
9200+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign(__ubuf__ fp8_e8m0_t* src)
9201+{
9202+ vector_fp8_e8m0_t dst;
9203+ asc_loadalign_impl(dst, src);
9204+ return dst;
9205+}
9206+ 
9207+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign(__ubuf__ fp8_e5m2_t* src)
9208+{
9209+ vector_fp8_e5m2_t dst;
9210+ asc_loadalign_impl(dst, src);
9211+ return dst;
9212+}
9213+ 
9214+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign(__ubuf__ fp8_e4m3fn_t* src)
9215+{
9216+ vector_fp8_e4m3fn_t dst;
9217+ asc_loadalign_impl(dst, src);
9218+ return dst;
9219+}
9220+ 
9221+__simd_callee__ inline vector_int16_t asc_loadalign(__ubuf__ int16_t* src)
9222+{
9223+ vector_int16_t dst;
9224+ asc_loadalign_impl(dst, src);
9225+ return dst;
9226+}
9227+ 
9228+__simd_callee__ inline vector_uint16_t asc_loadalign(__ubuf__ uint16_t* src)
9229+{
9230+ vector_uint16_t dst;
9231+ asc_loadalign_impl(dst, src);
9232+ return dst;
9233+}
9234+ 
9235+__simd_callee__ inline vector_half asc_loadalign(__ubuf__ half* src)
9236+{
9237+ vector_half dst;
9238+ asc_loadalign_impl(dst, src);
9239+ return dst;
9240+}
9241+ 
9242+__simd_callee__ inline vector_bfloat16_t asc_loadalign(__ubuf__ bfloat16_t* src)
9243+{
9244+ vector_bfloat16_t dst;
9245+ asc_loadalign_impl(dst, src);
9246+ return dst;
9247+}
9248+ 
9249+__simd_callee__ inline vector_int32_t asc_loadalign(__ubuf__ int32_t* src)
9250+{
9251+ vector_int32_t dst;
9252+ asc_loadalign_impl(dst, src);
9253+ return dst;
9254+}
9255+ 
9256+__simd_callee__ inline vector_uint32_t asc_loadalign(__ubuf__ uint32_t* src)
9257+{
9258+ vector_uint32_t dst;
9259+ asc_loadalign_impl(dst, src);
9260+ return dst;
9261+}
9262+ 
9263+__simd_callee__ inline vector_float asc_loadalign(__ubuf__ float* src)
9264+{
9265+ vector_float dst;
9266+ asc_loadalign_impl(dst, src);
9267+ return dst;
9268+}
9269+ 
9270+// ========== asc_loadalign_brc_datablock ==========
9271+__simd_callee__ inline vector_int4x2_t asc_loadalign_brc_datablock(__ubuf__ int4b_t* src)
9272+{
9273+ vector_int4x2_t dst;
9274+ asc_loadalign_brc_datablock_impl(dst, src);
9275+ return dst;
9276+}
9277+ 
9278+__simd_callee__ inline vector_int8_t asc_loadalign_brc_datablock(__ubuf__ int8_t* src)
9279+{
9280+ vector_int8_t dst;
9281+ asc_loadalign_brc_datablock_impl(dst, src);
9282+ return dst;
9283+}
9284+ 
9285+__simd_callee__ inline vector_uint8_t asc_loadalign_brc_datablock(__ubuf__ uint8_t* src)
9286+{
9287+ vector_uint8_t dst;
9288+ asc_loadalign_brc_datablock_impl(dst, src);
9289+ return dst;
9290+}
9291+ 
9292+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_brc_datablock(__ubuf__ fp4x2_e2m1_t* src)
9293+{
9294+ vector_fp4x2_e2m1_t dst;
9295+ asc_loadalign_brc_datablock_impl(dst, src);
9296+ return dst;
9297+}
9298+ 
9299+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_brc_datablock(__ubuf__ fp4x2_e1m2_t* src)
9300+{
9301+ vector_fp4x2_e1m2_t dst;
9302+ asc_loadalign_brc_datablock_impl(dst, src);
9303+ return dst;
9304+}
9305+ 
9306+__simd_callee__ inline vector_hifloat8_t asc_loadalign_brc_datablock(__ubuf__ hifloat8_t* src)
9307+{
9308+ vector_hifloat8_t dst;
9309+ asc_loadalign_brc_datablock_impl(dst, src);
9310+ return dst;
9311+}
9312+ 
9313+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_brc_datablock(__ubuf__ fp8_e8m0_t* src)
9314+{
9315+ vector_fp8_e8m0_t dst;
9316+ asc_loadalign_brc_datablock_impl(dst, src);
9317+ return dst;
9318+}
9319+ 
9320+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_brc_datablock(__ubuf__ fp8_e5m2_t* src)
9321+{
9322+ vector_fp8_e5m2_t dst;
9323+ asc_loadalign_brc_datablock_impl(dst, src);
9324+ return dst;
9325+}
9326+ 
9327+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_brc_datablock(__ubuf__ fp8_e4m3fn_t* src)
9328+{
9329+ vector_fp8_e4m3fn_t dst;
9330+ asc_loadalign_brc_datablock_impl(dst, src);
9331+ return dst;
9332+}
9333+ 
9334+__simd_callee__ inline vector_int16_t asc_loadalign_brc_datablock(__ubuf__ int16_t* src)
9335+{
9336+ vector_int16_t dst;
9337+ asc_loadalign_brc_datablock_impl(dst, src);
9338+ return dst;
9339+}
9340+ 
9341+__simd_callee__ inline vector_uint16_t asc_loadalign_brc_datablock(__ubuf__ uint16_t* src)
9342+{
9343+ vector_uint16_t dst;
9344+ asc_loadalign_brc_datablock_impl(dst, src);
9345+ return dst;
9346+}
9347+ 
9348+__simd_callee__ inline vector_half asc_loadalign_brc_datablock(__ubuf__ half* src)
9349+{
9350+ vector_half dst;
9351+ asc_loadalign_brc_datablock_impl(dst, src);
9352+ return dst;
9353+}
9354+ 
9355+__simd_callee__ inline vector_bfloat16_t asc_loadalign_brc_datablock(__ubuf__ bfloat16_t* src)
9356+{
9357+ vector_bfloat16_t dst;
9358+ asc_loadalign_brc_datablock_impl(dst, src);
9359+ return dst;
9360+}
9361+ 
9362+__simd_callee__ inline vector_int32_t asc_loadalign_brc_datablock(__ubuf__ int32_t* src)
9363+{
9364+ vector_int32_t dst;
9365+ asc_loadalign_brc_datablock_impl(dst, src);
9366+ return dst;
9367+}
9368+ 
9369+__simd_callee__ inline vector_uint32_t asc_loadalign_brc_datablock(__ubuf__ uint32_t* src)
9370+{
9371+ vector_uint32_t dst;
9372+ asc_loadalign_brc_datablock_impl(dst, src);
9373+ return dst;
9374+}
9375+ 
9376+__simd_callee__ inline vector_float asc_loadalign_brc_datablock(__ubuf__ float* src)
9377+{
9378+ vector_float dst;
9379+ asc_loadalign_brc_datablock_impl(dst, src);
9380+ return dst;
9381+}
9382+ 
9383+// ========== asc_loadalign_brc_elem ==========
9384+__simd_callee__ inline vector_int4x2_t asc_loadalign_brc_elem(__ubuf__ int4b_t* src)
9385+{
9386+ vector_int4x2_t dst;
9387+ asc_loadalign_brc_elem_impl(dst, src);
9388+ return dst;
9389+}
9390+ 
9391+__simd_callee__ inline vector_int8_t asc_loadalign_brc_elem(__ubuf__ int8_t* src)
9392+{
9393+ vector_int8_t dst;
9394+ asc_loadalign_brc_elem_impl(dst, src);
9395+ return dst;
9396+}
9397+ 
9398+__simd_callee__ inline vector_uint8_t asc_loadalign_brc_elem(__ubuf__ uint8_t* src)
9399+{
9400+ vector_uint8_t dst;
9401+ asc_loadalign_brc_elem_impl(dst, src);
9402+ return dst;
9403+}
9404+ 
9405+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_brc_elem(__ubuf__ fp4x2_e2m1_t* src)
9406+{
9407+ vector_fp4x2_e2m1_t dst;
9408+ asc_loadalign_brc_elem_impl(dst, src);
9409+ return dst;
9410+}
9411+ 
9412+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_brc_elem(__ubuf__ fp4x2_e1m2_t* src)
9413+{
9414+ vector_fp4x2_e1m2_t dst;
9415+ asc_loadalign_brc_elem_impl(dst, src);
9416+ return dst;
9417+}
9418+ 
9419+__simd_callee__ inline vector_hifloat8_t asc_loadalign_brc_elem(__ubuf__ hifloat8_t* src)
9420+{
9421+ vector_hifloat8_t dst;
9422+ asc_loadalign_brc_elem_impl(dst, src);
9423+ return dst;
9424+}
9425+ 
9426+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_brc_elem(__ubuf__ fp8_e8m0_t* src)
9427+{
9428+ vector_fp8_e8m0_t dst;
9429+ asc_loadalign_brc_elem_impl(dst, src);
9430+ return dst;
9431+}
9432+ 
9433+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_brc_elem(__ubuf__ fp8_e5m2_t* src)
9434+{
9435+ vector_fp8_e5m2_t dst;
9436+ asc_loadalign_brc_elem_impl(dst, src);
9437+ return dst;
9438+}
9439+ 
9440+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_brc_elem(__ubuf__ fp8_e4m3fn_t* src)
9441+{
9442+ vector_fp8_e4m3fn_t dst;
9443+ asc_loadalign_brc_elem_impl(dst, src);
9444+ return dst;
9445+}
9446+ 
9447+__simd_callee__ inline vector_int16_t asc_loadalign_brc_elem(__ubuf__ int16_t* src)
9448+{
9449+ vector_int16_t dst;
9450+ asc_loadalign_brc_elem_impl(dst, src);
9451+ return dst;
9452+}
9453+ 
9454+__simd_callee__ inline vector_uint16_t asc_loadalign_brc_elem(__ubuf__ uint16_t* src)
9455+{
9456+ vector_uint16_t dst;
9457+ asc_loadalign_brc_elem_impl(dst, src);
9458+ return dst;
9459+}
9460+ 
9461+__simd_callee__ inline vector_half asc_loadalign_brc_elem(__ubuf__ half* src)
9462+{
9463+ vector_half dst;
9464+ asc_loadalign_brc_elem_impl(dst, src);
9465+ return dst;
9466+}
9467+ 
9468+__simd_callee__ inline vector_bfloat16_t asc_loadalign_brc_elem(__ubuf__ bfloat16_t* src)
9469+{
9470+ vector_bfloat16_t dst;
9471+ asc_loadalign_brc_elem_impl(dst, src);
9472+ return dst;
9473+}
9474+ 
9475+__simd_callee__ inline vector_int32_t asc_loadalign_brc_elem(__ubuf__ int32_t* src)
9476+{
9477+ vector_int32_t dst;
9478+ asc_loadalign_brc_elem_impl(dst, src);
9479+ return dst;
9480+}
9481+ 
9482+__simd_callee__ inline vector_uint32_t asc_loadalign_brc_elem(__ubuf__ uint32_t* src)
9483+{
9484+ vector_uint32_t dst;
9485+ asc_loadalign_brc_elem_impl(dst, src);
9486+ return dst;
9487+}
9488+ 
9489+__simd_callee__ inline vector_float asc_loadalign_brc_elem(__ubuf__ float* src)
9490+{
9491+ vector_float dst;
9492+ asc_loadalign_brc_elem_impl(dst, src);
9493+ return dst;
9494+}
9495+ 
9496+// ========== asc_loadalign_brc_elem2datablock ==========
9497+__simd_callee__ inline vector_int16_t asc_loadalign_brc_elem2datablock(__ubuf__ int16_t* src)
9498+{
9499+ vector_int16_t dst;
9500+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9501+ return dst;
9502+}
9503+ 
9504+__simd_callee__ inline vector_uint16_t asc_loadalign_brc_elem2datablock(__ubuf__ uint16_t* src)
9505+{
9506+ vector_uint16_t dst;
9507+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9508+ return dst;
9509+}
9510+ 
9511+__simd_callee__ inline vector_half asc_loadalign_brc_elem2datablock(__ubuf__ half* src)
9512+{
9513+ vector_half dst;
9514+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9515+ return dst;
9516+}
9517+ 
9518+__simd_callee__ inline vector_bfloat16_t asc_loadalign_brc_elem2datablock(__ubuf__ bfloat16_t* src)
9519+{
9520+ vector_bfloat16_t dst;
9521+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9522+ return dst;
9523+}
9524+ 
9525+__simd_callee__ inline vector_int32_t asc_loadalign_brc_elem2datablock(__ubuf__ int32_t* src)
9526+{
9527+ vector_int32_t dst;
9528+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9529+ return dst;
9530+}
9531+ 
9532+__simd_callee__ inline vector_uint32_t asc_loadalign_brc_elem2datablock(__ubuf__ uint32_t* src)
9533+{
9534+ vector_uint32_t dst;
9535+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9536+ return dst;
9537+}
9538+ 
9539+__simd_callee__ inline vector_float asc_loadalign_brc_elem2datablock(__ubuf__ float* src)
9540+{
9541+ vector_float dst;
9542+ asc_loadalign_brc_elem2datablock_impl(dst, src);
9543+ return dst;
9544+}
9545+ 
9546+// ========== asc_loadalign_downsample ==========
9547+__simd_callee__ inline vector_int4x2_t asc_loadalign_downsample(__ubuf__ int4b_t* src)
9548+{
9549+ vector_int4x2_t dst;
9550+ asc_loadalign_downsample_impl(dst, src);
9551+ return dst;
9552+}
9553+ 
9554+__simd_callee__ inline vector_int8_t asc_loadalign_downsample(__ubuf__ int8_t* src)
9555+{
9556+ vector_int8_t dst;
9557+ asc_loadalign_downsample_impl(dst, src);
9558+ return dst;
9559+}
9560+ 
9561+__simd_callee__ inline vector_uint8_t asc_loadalign_downsample(__ubuf__ uint8_t* src)
9562+{
9563+ vector_uint8_t dst;
9564+ asc_loadalign_downsample_impl(dst, src);
9565+ return dst;
9566+}
9567+ 
9568+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_downsample(__ubuf__ fp4x2_e2m1_t* src)
9569+{
9570+ vector_fp4x2_e2m1_t dst;
9571+ asc_loadalign_downsample_impl(dst, src);
9572+ return dst;
9573+}
9574+ 
9575+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_downsample(__ubuf__ fp4x2_e1m2_t* src)
9576+{
9577+ vector_fp4x2_e1m2_t dst;
9578+ asc_loadalign_downsample_impl(dst, src);
9579+ return dst;
9580+}
9581+ 
9582+__simd_callee__ inline vector_hifloat8_t asc_loadalign_downsample(__ubuf__ hifloat8_t* src)
9583+{
9584+ vector_hifloat8_t dst;
9585+ asc_loadalign_downsample_impl(dst, src);
9586+ return dst;
9587+}
9588+ 
9589+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_downsample(__ubuf__ fp8_e8m0_t* src)
9590+{
9591+ vector_fp8_e8m0_t dst;
9592+ asc_loadalign_downsample_impl(dst, src);
9593+ return dst;
9594+}
9595+ 
9596+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_downsample(__ubuf__ fp8_e5m2_t* src)
9597+{
9598+ vector_fp8_e5m2_t dst;
9599+ asc_loadalign_downsample_impl(dst, src);
9600+ return dst;
9601+}
9602+ 
9603+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_downsample(__ubuf__ fp8_e4m3fn_t* src)
9604+{
9605+ vector_fp8_e4m3fn_t dst;
9606+ asc_loadalign_downsample_impl(dst, src);
9607+ return dst;
9608+}
9609+ 
9610+__simd_callee__ inline vector_int16_t asc_loadalign_downsample(__ubuf__ int16_t* src)
9611+{
9612+ vector_int16_t dst;
9613+ asc_loadalign_downsample_impl(dst, src);
9614+ return dst;
9615+}
9616+ 
9617+__simd_callee__ inline vector_uint16_t asc_loadalign_downsample(__ubuf__ uint16_t* src)
9618+{
9619+ vector_uint16_t dst;
9620+ asc_loadalign_downsample_impl(dst, src);
9621+ return dst;
9622+}
9623+ 
9624+__simd_callee__ inline vector_half asc_loadalign_downsample(__ubuf__ half* src)
9625+{
9626+ vector_half dst;
9627+ asc_loadalign_downsample_impl(dst, src);
9628+ return dst;
9629+}
9630+ 
9631+__simd_callee__ inline vector_bfloat16_t asc_loadalign_downsample(__ubuf__ bfloat16_t* src)
9632+{
9633+ vector_bfloat16_t dst;
9634+ asc_loadalign_downsample_impl(dst, src);
9635+ return dst;
9636+}
9637+ 
9638+// ========== asc_loadalign_unpack ==========
9639+__simd_callee__ inline vector_int4x2_t asc_loadalign_unpack(__ubuf__ int4b_t* src)
9640+{
9641+ vector_int4x2_t dst;
9642+ asc_loadalign_unpack_impl(dst, src);
9643+ return dst;
9644+}
9645+ 
9646+__simd_callee__ inline vector_int8_t asc_loadalign_unpack(__ubuf__ int8_t* src)
9647+{
9648+ vector_int8_t dst;
9649+ asc_loadalign_unpack_impl(dst, src);
9650+ return dst;
9651+}
9652+ 
9653+__simd_callee__ inline vector_uint8_t asc_loadalign_unpack(__ubuf__ uint8_t* src)
9654+{
9655+ vector_uint8_t dst;
9656+ asc_loadalign_unpack_impl(dst, src);
9657+ return dst;
9658+}
9659+ 
9660+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_unpack(__ubuf__ fp4x2_e2m1_t* src)
9661+{
9662+ vector_fp4x2_e2m1_t dst;
9663+ asc_loadalign_unpack_impl(dst, src);
9664+ return dst;
9665+}
9666+ 
9667+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_unpack(__ubuf__ fp4x2_e1m2_t* src)
9668+{
9669+ vector_fp4x2_e1m2_t dst;
9670+ asc_loadalign_unpack_impl(dst, src);
9671+ return dst;
9672+}
9673+ 
9674+__simd_callee__ inline vector_hifloat8_t asc_loadalign_unpack(__ubuf__ hifloat8_t* src)
9675+{
9676+ vector_hifloat8_t dst;
9677+ asc_loadalign_unpack_impl(dst, src);
9678+ return dst;
9679+}
9680+ 
9681+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_unpack(__ubuf__ fp8_e8m0_t* src)
9682+{
9683+ vector_fp8_e8m0_t dst;
9684+ asc_loadalign_unpack_impl(dst, src);
9685+ return dst;
9686+}
9687+ 
9688+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_unpack(__ubuf__ fp8_e5m2_t* src)
9689+{
9690+ vector_fp8_e5m2_t dst;
9691+ asc_loadalign_unpack_impl(dst, src);
9692+ return dst;
9693+}
9694+ 
9695+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_unpack(__ubuf__ fp8_e4m3fn_t* src)
9696+{
9697+ vector_fp8_e4m3fn_t dst;
9698+ asc_loadalign_unpack_impl(dst, src);
9699+ return dst;
9700+}
9701+ 
9702+__simd_callee__ inline vector_int16_t asc_loadalign_unpack(__ubuf__ int16_t* src)
9703+{
9704+ vector_int16_t dst;
9705+ asc_loadalign_unpack_impl(dst, src);
9706+ return dst;
9707+}
9708+ 
9709+__simd_callee__ inline vector_uint16_t asc_loadalign_unpack(__ubuf__ uint16_t* src)
9710+{
9711+ vector_uint16_t dst;
9712+ asc_loadalign_unpack_impl(dst, src);
9713+ return dst;
9714+}
9715+ 
9716+__simd_callee__ inline vector_half asc_loadalign_unpack(__ubuf__ half* src)
9717+{
9718+ vector_half dst;
9719+ asc_loadalign_unpack_impl(dst, src);
9720+ return dst;
9721+}
9722+ 
9723+__simd_callee__ inline vector_bfloat16_t asc_loadalign_unpack(__ubuf__ bfloat16_t* src)
9724+{
9725+ vector_bfloat16_t dst;
9726+ asc_loadalign_unpack_impl(dst, src);
9727+ return dst;
9728+}
9729+ 
9730+__simd_callee__ inline vector_int32_t asc_loadalign_unpack(__ubuf__ int32_t* src)
9731+{
9732+ vector_int32_t dst;
9733+ asc_loadalign_unpack_impl(dst, src);
9734+ return dst;
9735+}
9736+ 
9737+__simd_callee__ inline vector_uint32_t asc_loadalign_unpack(__ubuf__ uint32_t* src)
9738+{
9739+ vector_uint32_t dst;
9740+ asc_loadalign_unpack_impl(dst, src);
9741+ return dst;
9742+}
9743+ 
9744+__simd_callee__ inline vector_float asc_loadalign_unpack(__ubuf__ float* src)
9745+{
9746+ vector_float dst;
9747+ asc_loadalign_unpack_impl(dst, src);
9748+ return dst;
9749+}
9750+ 
9751+// ========== asc_loadalign_unpack4 ==========
9752+__simd_callee__ inline vector_int4x2_t asc_loadalign_unpack4(__ubuf__ int4b_t* src)
9753+{
9754+ vector_int4x2_t dst;
9755+ asc_loadalign_unpack4_impl(dst, src);
9756+ return dst;
9757+}
9758+ 
9759+__simd_callee__ inline vector_int8_t asc_loadalign_unpack4(__ubuf__ int8_t* src)
9760+{
9761+ vector_int8_t dst;
9762+ asc_loadalign_unpack4_impl(dst, src);
9763+ return dst;
9764+}
9765+ 
9766+__simd_callee__ inline vector_uint8_t asc_loadalign_unpack4(__ubuf__ uint8_t* src)
9767+{
9768+ vector_uint8_t dst;
9769+ asc_loadalign_unpack4_impl(dst, src);
9770+ return dst;
9771+}
9772+ 
9773+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_unpack4(__ubuf__ fp4x2_e2m1_t* src)
9774+{
9775+ vector_fp4x2_e2m1_t dst;
9776+ asc_loadalign_unpack4_impl(dst, src);
9777+ return dst;
9778+}
9779+ 
9780+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_unpack4(__ubuf__ fp4x2_e1m2_t* src)
9781+{
9782+ vector_fp4x2_e1m2_t dst;
9783+ asc_loadalign_unpack4_impl(dst, src);
9784+ return dst;
9785+}
9786+ 
9787+__simd_callee__ inline vector_hifloat8_t asc_loadalign_unpack4(__ubuf__ hifloat8_t* src)
9788+{
9789+ vector_hifloat8_t dst;
9790+ asc_loadalign_unpack4_impl(dst, src);
9791+ return dst;
9792+}
9793+ 
9794+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_unpack4(__ubuf__ fp8_e8m0_t* src)
9795+{
9796+ vector_fp8_e8m0_t dst;
9797+ asc_loadalign_unpack4_impl(dst, src);
9798+ return dst;
9799+}
9800+ 
9801+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_unpack4(__ubuf__ fp8_e5m2_t* src)
9802+{
9803+ vector_fp8_e5m2_t dst;
9804+ asc_loadalign_unpack4_impl(dst, src);
9805+ return dst;
9806+}
9807+ 
9808+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_unpack4(__ubuf__ fp8_e4m3fn_t* src)
9809+{
9810+ vector_fp8_e4m3fn_t dst;
9811+ asc_loadalign_unpack4_impl(dst, src);
9812+ return dst;
9813+}
9814+ 
9815+// ========== asc_loadalign_upsample ==========
9816+__simd_callee__ inline vector_int4x2_t asc_loadalign_upsample(__ubuf__ int4b_t* src)
9817+{
9818+ vector_int4x2_t dst;
9819+ asc_loadalign_upsample_impl(dst, src);
9820+ return dst;
9821+}
9822+ 
9823+__simd_callee__ inline vector_int8_t asc_loadalign_upsample(__ubuf__ int8_t* src)
9824+{
9825+ vector_int8_t dst;
9826+ asc_loadalign_upsample_impl(dst, src);
9827+ return dst;
9828+}
9829+ 
9830+__simd_callee__ inline vector_uint8_t asc_loadalign_upsample(__ubuf__ uint8_t* src)
9831+{
9832+ vector_uint8_t dst;
9833+ asc_loadalign_upsample_impl(dst, src);
9834+ return dst;
9835+}
9836+ 
9837+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_upsample(__ubuf__ fp4x2_e2m1_t* src)
9838+{
9839+ vector_fp4x2_e2m1_t dst;
9840+ asc_loadalign_upsample_impl(dst, src);
9841+ return dst;
9842+}
9843+ 
9844+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_upsample(__ubuf__ fp4x2_e1m2_t* src)
9845+{
9846+ vector_fp4x2_e1m2_t dst;
9847+ asc_loadalign_upsample_impl(dst, src);
9848+ return dst;
9849+}
9850+ 
9851+__simd_callee__ inline vector_hifloat8_t asc_loadalign_upsample(__ubuf__ hifloat8_t* src)
9852+{
9853+ vector_hifloat8_t dst;
9854+ asc_loadalign_upsample_impl(dst, src);
9855+ return dst;
9856+}
9857+ 
9858+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_upsample(__ubuf__ fp8_e8m0_t* src)
9859+{
9860+ vector_fp8_e8m0_t dst;
9861+ asc_loadalign_upsample_impl(dst, src);
9862+ return dst;
9863+}
9864+ 
9865+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_upsample(__ubuf__ fp8_e5m2_t* src)
9866+{
9867+ vector_fp8_e5m2_t dst;
9868+ asc_loadalign_upsample_impl(dst, src);
9869+ return dst;
9870+}
9871+ 
9872+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_upsample(__ubuf__ fp8_e4m3fn_t* src)
9873+{
9874+ vector_fp8_e4m3fn_t dst;
9875+ asc_loadalign_upsample_impl(dst, src);
9876+ return dst;
9877+}
9878+ 
9879+__simd_callee__ inline vector_int16_t asc_loadalign_upsample(__ubuf__ int16_t* src)
9880+{
9881+ vector_int16_t dst;
9882+ asc_loadalign_upsample_impl(dst, src);
9883+ return dst;
9884+}
9885+ 
9886+__simd_callee__ inline vector_uint16_t asc_loadalign_upsample(__ubuf__ uint16_t* src)
9887+{
9888+ vector_uint16_t dst;
9889+ asc_loadalign_upsample_impl(dst, src);
9890+ return dst;
9891+}
9892+ 
9893+__simd_callee__ inline vector_half asc_loadalign_upsample(__ubuf__ half* src)
9894+{
9895+ vector_half dst;
9896+ asc_loadalign_upsample_impl(dst, src);
9897+ return dst;
9898+}
9899+ 
9900+__simd_callee__ inline vector_bfloat16_t asc_loadalign_upsample(__ubuf__ bfloat16_t* src)
9901+{
9902+ vector_bfloat16_t dst;
9903+ asc_loadalign_upsample_impl(dst, src);
9904+ return dst;
9905+}
9906+ 
9907+// ========== asc_loadalign_datablock_strided ==========
9908+__simd_callee__ inline vector_int4x2_t asc_loadalign_datablock_strided(
9909+ __ubuf__ int4b_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9910+{
9911+ vector_int4x2_t dst;
9912+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9913+ return dst;
9914+}
9915+ 
9916+__simd_callee__ inline vector_int8_t asc_loadalign_datablock_strided(
9917+ __ubuf__ int8_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9918+{
9919+ vector_int8_t dst;
9920+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9921+ return dst;
9922+}
9923+ 
9924+__simd_callee__ inline vector_uint8_t asc_loadalign_datablock_strided(
9925+ __ubuf__ uint8_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9926+{
9927+ vector_uint8_t dst;
9928+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9929+ return dst;
9930+}
9931+ 
9932+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_datablock_strided(
9933+ __ubuf__ fp4x2_e2m1_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9934+{
9935+ vector_fp4x2_e2m1_t dst;
9936+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9937+ return dst;
9938+}
9939+ 
9940+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_datablock_strided(
9941+ __ubuf__ fp4x2_e1m2_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9942+{
9943+ vector_fp4x2_e1m2_t dst;
9944+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9945+ return dst;
9946+}
9947+ 
9948+__simd_callee__ inline vector_hifloat8_t asc_loadalign_datablock_strided(
9949+ __ubuf__ hifloat8_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9950+{
9951+ vector_hifloat8_t dst;
9952+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9953+ return dst;
9954+}
9955+ 
9956+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_datablock_strided(
9957+ __ubuf__ fp8_e8m0_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9958+{
9959+ vector_fp8_e8m0_t dst;
9960+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9961+ return dst;
9962+}
9963+ 
9964+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_datablock_strided(
9965+ __ubuf__ fp8_e5m2_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9966+{
9967+ vector_fp8_e5m2_t dst;
9968+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9969+ return dst;
9970+}
9971+ 
9972+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_datablock_strided(
9973+ __ubuf__ fp8_e4m3fn_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9974+{
9975+ vector_fp8_e4m3fn_t dst;
9976+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9977+ return dst;
9978+}
9979+ 
9980+__simd_callee__ inline vector_int16_t asc_loadalign_datablock_strided(
9981+ __ubuf__ int16_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9982+{
9983+ vector_int16_t dst;
9984+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9985+ return dst;
9986+}
9987+ 
9988+__simd_callee__ inline vector_uint16_t asc_loadalign_datablock_strided(
9989+ __ubuf__ uint16_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9990+{
9991+ vector_uint16_t dst;
9992+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
9993+ return dst;
9994+}
9995+ 
9996+__simd_callee__ inline vector_half asc_loadalign_datablock_strided(
9997+ __ubuf__ half* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
9998+{
9999+ vector_half dst;
10000+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
10001+ return dst;
10002+}
10003+ 
10004+__simd_callee__ inline vector_bfloat16_t asc_loadalign_datablock_strided(
10005+ __ubuf__ bfloat16_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
10006+{
10007+ vector_bfloat16_t dst;
10008+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
10009+ return dst;
10010+}
10011+ 
10012+__simd_callee__ inline vector_int32_t asc_loadalign_datablock_strided(
10013+ __ubuf__ int32_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
10014+{
10015+ vector_int32_t dst;
10016+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
10017+ return dst;
10018+}
10019+ 
10020+__simd_callee__ inline vector_uint32_t asc_loadalign_datablock_strided(
10021+ __ubuf__ uint32_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
10022+{
10023+ vector_uint32_t dst;
10024+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
10025+ return dst;
10026+}
10027+ 
10028+__simd_callee__ inline vector_float asc_loadalign_datablock_strided(
10029+ __ubuf__ float* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask)
10030+{
10031+ vector_float dst;
10032+ asc_loadalign_impl(dst, src, block_stride, repeat_stride, mask);
10033+ return dst;
10034+}
10035+ 
10036+// ========== return mask register APIs ==========
10037+__simd_callee__ inline vector_bool asc_loadalign_mask(__ubuf__ uint32_t* src)
10038+{
10039+ vector_bool dst;
10040+ asc_loadalign_impl(dst, src);
10041+ return dst;
10042+}
10043+ 
10044+__simd_callee__ inline vector_bool asc_loadalign_mask_downsample(__ubuf__ uint32_t* src)
10045+{
10046+ vector_bool dst;
10047+ asc_loadalign_downsample_impl(dst, src);
10048+ return dst;
10049+}
10050+ 
10051+__simd_callee__ inline vector_bool asc_loadalign_mask_upsample(__ubuf__ uint32_t* src)
10052+{
10053+ vector_bool dst;
10054+ asc_loadalign_upsample_impl(dst, src);
10055+ return dst;
10056+}
10057+ 
9042// ========== asc_storealign_pack_quarter ==========10058// ========== asc_storealign_pack_quarter ==========
9043__simd_callee__ inline void asc_storealign_pack_quarter(10059__simd_callee__ inline void asc_storealign_pack_quarter(
9044 __ubuf__ int32_t* dst_align32b, vector_int32_t src, vector_bool mask)10060 __ubuf__ int32_t* dst_align32b, vector_int32_t src, vector_bool mask)
Minclude/c_api/composite/loadunalign_composite.h+34-0
@@ -25,6 +25,40 @@
25#include "impl/c_api/instr_impl/npu_arch_3510/vector_datamove_impl.h"25#include "impl/c_api/instr_impl/npu_arch_3510/vector_datamove_impl.h"
26#endif26#endif
27 27 
28+// ========== return-value load APIs ==========
29+// ========== asc_load ==========
30+__simd_callee__ inline vector_int4x2_t asc_load(__ubuf__ int4b_t* src);
31+ 
32+__simd_callee__ inline vector_int8_t asc_load(__ubuf__ int8_t* src);
33+ 
34+__simd_callee__ inline vector_uint8_t asc_load(__ubuf__ uint8_t* src);
35+ 
36+__simd_callee__ inline vector_fp4x2_e2m1_t asc_load(__ubuf__ fp4x2_e2m1_t* src);
37+ 
38+__simd_callee__ inline vector_fp4x2_e1m2_t asc_load(__ubuf__ fp4x2_e1m2_t* src);
39+ 
40+__simd_callee__ inline vector_hifloat8_t asc_load(__ubuf__ hifloat8_t* src);
41+ 
42+__simd_callee__ inline vector_fp8_e8m0_t asc_load(__ubuf__ fp8_e8m0_t* src);
43+ 
44+__simd_callee__ inline vector_fp8_e5m2_t asc_load(__ubuf__ fp8_e5m2_t* src);
45+ 
46+__simd_callee__ inline vector_fp8_e4m3fn_t asc_load(__ubuf__ fp8_e4m3fn_t* src);
47+ 
48+__simd_callee__ inline vector_int16_t asc_load(__ubuf__ int16_t* src);
49+ 
50+__simd_callee__ inline vector_uint16_t asc_load(__ubuf__ uint16_t* src);
51+ 
52+__simd_callee__ inline vector_half asc_load(__ubuf__ half* src);
53+ 
54+__simd_callee__ inline vector_bfloat16_t asc_load(__ubuf__ bfloat16_t* src);
55+ 
56+__simd_callee__ inline vector_int32_t asc_load(__ubuf__ int32_t* src);
57+ 
58+__simd_callee__ inline vector_uint32_t asc_load(__ubuf__ uint32_t* src);
59+ 
60+__simd_callee__ inline vector_float asc_load(__ubuf__ float* src);
61+ 
28__simd_callee__ inline void asc_load(vector_int8_t& dst, __ubuf__ int8_t* src);62__simd_callee__ inline void asc_load(vector_int8_t& dst, __ubuf__ int8_t* src);
29 63 
30__simd_callee__ inline void asc_load(vector_uint8_t& dst, __ubuf__ uint8_t* src);64__simd_callee__ inline void asc_load(vector_uint8_t& dst, __ubuf__ uint8_t* src);
Minclude/c_api/reg_compute/load/loadalign.h+278-0
@@ -25,6 +25,284 @@
25#include "impl/c_api/instr_impl/npu_arch_3510/vector_datamove_impl.h"25#include "impl/c_api/instr_impl/npu_arch_3510/vector_datamove_impl.h"
26#endif26#endif
27 27 
28+// ========== return-value load APIs ==========
29+// ========== return vector data register APIs ==========
30+// ========== asc_loadalign ==========
31+__simd_callee__ inline vector_int4x2_t asc_loadalign(__ubuf__ int4b_t* src);
32+ 
33+__simd_callee__ inline vector_int8_t asc_loadalign(__ubuf__ int8_t* src);
34+ 
35+__simd_callee__ inline vector_uint8_t asc_loadalign(__ubuf__ uint8_t* src);
36+ 
37+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign(__ubuf__ fp4x2_e2m1_t* src);
38+ 
39+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign(__ubuf__ fp4x2_e1m2_t* src);
40+ 
41+__simd_callee__ inline vector_hifloat8_t asc_loadalign(__ubuf__ hifloat8_t* src);
42+ 
43+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign(__ubuf__ fp8_e8m0_t* src);
44+ 
45+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign(__ubuf__ fp8_e5m2_t* src);
46+ 
47+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign(__ubuf__ fp8_e4m3fn_t* src);
48+ 
49+__simd_callee__ inline vector_int16_t asc_loadalign(__ubuf__ int16_t* src);
50+ 
51+__simd_callee__ inline vector_uint16_t asc_loadalign(__ubuf__ uint16_t* src);
52+ 
53+__simd_callee__ inline vector_half asc_loadalign(__ubuf__ half* src);
54+ 
55+__simd_callee__ inline vector_bfloat16_t asc_loadalign(__ubuf__ bfloat16_t* src);
56+ 
57+__simd_callee__ inline vector_int32_t asc_loadalign(__ubuf__ int32_t* src);
58+ 
59+__simd_callee__ inline vector_uint32_t asc_loadalign(__ubuf__ uint32_t* src);
60+ 
61+__simd_callee__ inline vector_float asc_loadalign(__ubuf__ float* src);
62+ 
63+// ========== asc_loadalign_brc_datablock ==========
64+__simd_callee__ inline vector_int4x2_t asc_loadalign_brc_datablock(__ubuf__ int4b_t* src);
65+ 
66+__simd_callee__ inline vector_int8_t asc_loadalign_brc_datablock(__ubuf__ int8_t* src);
67+ 
68+__simd_callee__ inline vector_uint8_t asc_loadalign_brc_datablock(__ubuf__ uint8_t* src);
69+ 
70+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_brc_datablock(__ubuf__ fp4x2_e2m1_t* src);
71+ 
72+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_brc_datablock(__ubuf__ fp4x2_e1m2_t* src);
73+ 
74+__simd_callee__ inline vector_hifloat8_t asc_loadalign_brc_datablock(__ubuf__ hifloat8_t* src);
75+ 
76+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_brc_datablock(__ubuf__ fp8_e8m0_t* src);
77+ 
78+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_brc_datablock(__ubuf__ fp8_e5m2_t* src);
79+ 
80+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_brc_datablock(__ubuf__ fp8_e4m3fn_t* src);
81+ 
82+__simd_callee__ inline vector_int16_t asc_loadalign_brc_datablock(__ubuf__ int16_t* src);
83+ 
84+__simd_callee__ inline vector_uint16_t asc_loadalign_brc_datablock(__ubuf__ uint16_t* src);
85+ 
86+__simd_callee__ inline vector_half asc_loadalign_brc_datablock(__ubuf__ half* src);
87+ 
88+__simd_callee__ inline vector_bfloat16_t asc_loadalign_brc_datablock(__ubuf__ bfloat16_t* src);
89+ 
90+__simd_callee__ inline vector_int32_t asc_loadalign_brc_datablock(__ubuf__ int32_t* src);
91+ 
92+__simd_callee__ inline vector_uint32_t asc_loadalign_brc_datablock(__ubuf__ uint32_t* src);
93+ 
94+__simd_callee__ inline vector_float asc_loadalign_brc_datablock(__ubuf__ float* src);
95+ 
96+// ========== asc_loadalign_brc_elem ==========
97+__simd_callee__ inline vector_int4x2_t asc_loadalign_brc_elem(__ubuf__ int4b_t* src);
98+ 
99+__simd_callee__ inline vector_int8_t asc_loadalign_brc_elem(__ubuf__ int8_t* src);
100+ 
101+__simd_callee__ inline vector_uint8_t asc_loadalign_brc_elem(__ubuf__ uint8_t* src);
102+ 
103+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_brc_elem(__ubuf__ fp4x2_e2m1_t* src);
104+ 
105+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_brc_elem(__ubuf__ fp4x2_e1m2_t* src);
106+ 
107+__simd_callee__ inline vector_hifloat8_t asc_loadalign_brc_elem(__ubuf__ hifloat8_t* src);
108+ 
109+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_brc_elem(__ubuf__ fp8_e8m0_t* src);
110+ 
111+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_brc_elem(__ubuf__ fp8_e5m2_t* src);
112+ 
113+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_brc_elem(__ubuf__ fp8_e4m3fn_t* src);
114+ 
115+__simd_callee__ inline vector_int16_t asc_loadalign_brc_elem(__ubuf__ int16_t* src);
116+ 
117+__simd_callee__ inline vector_uint16_t asc_loadalign_brc_elem(__ubuf__ uint16_t* src);
118+ 
119+__simd_callee__ inline vector_half asc_loadalign_brc_elem(__ubuf__ half* src);
120+ 
121+__simd_callee__ inline vector_bfloat16_t asc_loadalign_brc_elem(__ubuf__ bfloat16_t* src);
122+ 
123+__simd_callee__ inline vector_int32_t asc_loadalign_brc_elem(__ubuf__ int32_t* src);
124+ 
125+__simd_callee__ inline vector_uint32_t asc_loadalign_brc_elem(__ubuf__ uint32_t* src);
126+ 
127+__simd_callee__ inline vector_float asc_loadalign_brc_elem(__ubuf__ float* src);
128+ 
129+// ========== asc_loadalign_brc_elem2datablock ==========
130+__simd_callee__ inline vector_int16_t asc_loadalign_brc_elem2datablock(__ubuf__ int16_t* src);
131+ 
132+__simd_callee__ inline vector_uint16_t asc_loadalign_brc_elem2datablock(__ubuf__ uint16_t* src);
133+ 
134+__simd_callee__ inline vector_half asc_loadalign_brc_elem2datablock(__ubuf__ half* src);
135+ 
136+__simd_callee__ inline vector_bfloat16_t asc_loadalign_brc_elem2datablock(__ubuf__ bfloat16_t* src);
137+ 
138+__simd_callee__ inline vector_int32_t asc_loadalign_brc_elem2datablock(__ubuf__ int32_t* src);
139+ 
140+__simd_callee__ inline vector_uint32_t asc_loadalign_brc_elem2datablock(__ubuf__ uint32_t* src);
141+ 
142+__simd_callee__ inline vector_float asc_loadalign_brc_elem2datablock(__ubuf__ float* src);
143+ 
144+// ========== asc_loadalign_downsample ==========
145+__simd_callee__ inline vector_int4x2_t asc_loadalign_downsample(__ubuf__ int4b_t* src);
146+ 
147+__simd_callee__ inline vector_int8_t asc_loadalign_downsample(__ubuf__ int8_t* src);
148+ 
149+__simd_callee__ inline vector_uint8_t asc_loadalign_downsample(__ubuf__ uint8_t* src);
150+ 
151+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_downsample(__ubuf__ fp4x2_e2m1_t* src);
152+ 
153+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_downsample(__ubuf__ fp4x2_e1m2_t* src);
154+ 
155+__simd_callee__ inline vector_hifloat8_t asc_loadalign_downsample(__ubuf__ hifloat8_t* src);
156+ 
157+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_downsample(__ubuf__ fp8_e8m0_t* src);
158+ 
159+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_downsample(__ubuf__ fp8_e5m2_t* src);
160+ 
161+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_downsample(__ubuf__ fp8_e4m3fn_t* src);
162+ 
163+__simd_callee__ inline vector_int16_t asc_loadalign_downsample(__ubuf__ int16_t* src);
164+ 
165+__simd_callee__ inline vector_uint16_t asc_loadalign_downsample(__ubuf__ uint16_t* src);
166+ 
167+__simd_callee__ inline vector_half asc_loadalign_downsample(__ubuf__ half* src);
168+ 
169+__simd_callee__ inline vector_bfloat16_t asc_loadalign_downsample(__ubuf__ bfloat16_t* src);
170+ 
171+// ========== asc_loadalign_unpack ==========
172+__simd_callee__ inline vector_int4x2_t asc_loadalign_unpack(__ubuf__ int4b_t* src);
173+ 
174+__simd_callee__ inline vector_int8_t asc_loadalign_unpack(__ubuf__ int8_t* src);
175+ 
176+__simd_callee__ inline vector_uint8_t asc_loadalign_unpack(__ubuf__ uint8_t* src);
177+ 
178+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_unpack(__ubuf__ fp4x2_e2m1_t* src);
179+ 
180+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_unpack(__ubuf__ fp4x2_e1m2_t* src);
181+ 
182+__simd_callee__ inline vector_hifloat8_t asc_loadalign_unpack(__ubuf__ hifloat8_t* src);
183+ 
184+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_unpack(__ubuf__ fp8_e8m0_t* src);
185+ 
186+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_unpack(__ubuf__ fp8_e5m2_t* src);
187+ 
188+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_unpack(__ubuf__ fp8_e4m3fn_t* src);
189+ 
190+__simd_callee__ inline vector_int16_t asc_loadalign_unpack(__ubuf__ int16_t* src);
191+ 
192+__simd_callee__ inline vector_uint16_t asc_loadalign_unpack(__ubuf__ uint16_t* src);
193+ 
194+__simd_callee__ inline vector_half asc_loadalign_unpack(__ubuf__ half* src);
195+ 
196+__simd_callee__ inline vector_bfloat16_t asc_loadalign_unpack(__ubuf__ bfloat16_t* src);
197+ 
198+__simd_callee__ inline vector_int32_t asc_loadalign_unpack(__ubuf__ int32_t* src);
199+ 
200+__simd_callee__ inline vector_uint32_t asc_loadalign_unpack(__ubuf__ uint32_t* src);
201+ 
202+__simd_callee__ inline vector_float asc_loadalign_unpack(__ubuf__ float* src);
203+ 
204+// ========== asc_loadalign_unpack4 ==========
205+__simd_callee__ inline vector_int4x2_t asc_loadalign_unpack4(__ubuf__ int4b_t* src);
206+ 
207+__simd_callee__ inline vector_int8_t asc_loadalign_unpack4(__ubuf__ int8_t* src);
208+ 
209+__simd_callee__ inline vector_uint8_t asc_loadalign_unpack4(__ubuf__ uint8_t* src);
210+ 
211+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_unpack4(__ubuf__ fp4x2_e2m1_t* src);
212+ 
213+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_unpack4(__ubuf__ fp4x2_e1m2_t* src);
214+ 
215+__simd_callee__ inline vector_hifloat8_t asc_loadalign_unpack4(__ubuf__ hifloat8_t* src);
216+ 
217+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_unpack4(__ubuf__ fp8_e8m0_t* src);
218+ 
219+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_unpack4(__ubuf__ fp8_e5m2_t* src);
220+ 
221+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_unpack4(__ubuf__ fp8_e4m3fn_t* src);
222+ 
223+// ========== asc_loadalign_upsample ==========
224+__simd_callee__ inline vector_int4x2_t asc_loadalign_upsample(__ubuf__ int4b_t* src);
225+ 
226+__simd_callee__ inline vector_int8_t asc_loadalign_upsample(__ubuf__ int8_t* src);
227+ 
228+__simd_callee__ inline vector_uint8_t asc_loadalign_upsample(__ubuf__ uint8_t* src);
229+ 
230+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_upsample(__ubuf__ fp4x2_e2m1_t* src);
231+ 
232+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_upsample(__ubuf__ fp4x2_e1m2_t* src);
233+ 
234+__simd_callee__ inline vector_hifloat8_t asc_loadalign_upsample(__ubuf__ hifloat8_t* src);
235+ 
236+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_upsample(__ubuf__ fp8_e8m0_t* src);
237+ 
238+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_upsample(__ubuf__ fp8_e5m2_t* src);
239+ 
240+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_upsample(__ubuf__ fp8_e4m3fn_t* src);
241+ 
242+__simd_callee__ inline vector_int16_t asc_loadalign_upsample(__ubuf__ int16_t* src);
243+ 
244+__simd_callee__ inline vector_uint16_t asc_loadalign_upsample(__ubuf__ uint16_t* src);
245+ 
246+__simd_callee__ inline vector_half asc_loadalign_upsample(__ubuf__ half* src);
247+ 
248+__simd_callee__ inline vector_bfloat16_t asc_loadalign_upsample(__ubuf__ bfloat16_t* src);
249+ 
250+// ========== asc_loadalign_datablock_strided ==========
251+__simd_callee__ inline vector_int4x2_t asc_loadalign_datablock_strided(
252+ __ubuf__ int4b_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
253+ 
254+__simd_callee__ inline vector_int8_t asc_loadalign_datablock_strided(
255+ __ubuf__ int8_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
256+ 
257+__simd_callee__ inline vector_uint8_t asc_loadalign_datablock_strided(
258+ __ubuf__ uint8_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
259+ 
260+__simd_callee__ inline vector_fp4x2_e2m1_t asc_loadalign_datablock_strided(
261+ __ubuf__ fp4x2_e2m1_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
262+ 
263+__simd_callee__ inline vector_fp4x2_e1m2_t asc_loadalign_datablock_strided(
264+ __ubuf__ fp4x2_e1m2_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
265+ 
266+__simd_callee__ inline vector_hifloat8_t asc_loadalign_datablock_strided(
267+ __ubuf__ hifloat8_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
268+ 
269+__simd_callee__ inline vector_fp8_e8m0_t asc_loadalign_datablock_strided(
270+ __ubuf__ fp8_e8m0_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
271+ 
272+__simd_callee__ inline vector_fp8_e5m2_t asc_loadalign_datablock_strided(
273+ __ubuf__ fp8_e5m2_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
274+ 
275+__simd_callee__ inline vector_fp8_e4m3fn_t asc_loadalign_datablock_strided(
276+ __ubuf__ fp8_e4m3fn_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
277+ 
278+__simd_callee__ inline vector_int16_t asc_loadalign_datablock_strided(
279+ __ubuf__ int16_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
280+ 
281+__simd_callee__ inline vector_uint16_t asc_loadalign_datablock_strided(
282+ __ubuf__ uint16_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
283+ 
284+__simd_callee__ inline vector_half asc_loadalign_datablock_strided(
285+ __ubuf__ half* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
286+ 
287+__simd_callee__ inline vector_bfloat16_t asc_loadalign_datablock_strided(
288+ __ubuf__ bfloat16_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
289+ 
290+__simd_callee__ inline vector_int32_t asc_loadalign_datablock_strided(
291+ __ubuf__ int32_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
292+ 
293+__simd_callee__ inline vector_uint32_t asc_loadalign_datablock_strided(
294+ __ubuf__ uint32_t* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
295+ 
296+__simd_callee__ inline vector_float asc_loadalign_datablock_strided(
297+ __ubuf__ float* src, uint16_t block_stride, uint16_t repeat_stride, vector_bool mask);
298+ 
299+// ========== return mask register APIs ==========
300+__simd_callee__ inline vector_bool asc_loadalign_mask(__ubuf__ uint32_t* src);
301+ 
302+__simd_callee__ inline vector_bool asc_loadalign_mask_downsample(__ubuf__ uint32_t* src);
303+ 
304+__simd_callee__ inline vector_bool asc_loadalign_mask_upsample(__ubuf__ uint32_t* src);
305+ 
28__simd_callee__ inline void asc_loadalign(vector_int8_t& dst, __ubuf__ int8_t* src);306__simd_callee__ inline void asc_loadalign(vector_int8_t& dst, __ubuf__ int8_t* src);
29 307 
30__simd_callee__ inline void asc_loadalign(vector_uint8_t& dst, __ubuf__ uint8_t* src);308__simd_callee__ inline void asc_loadalign(vector_uint8_t& dst, __ubuf__ uint8_t* src);
Mtests/api/c_api/ascendc_host_header_checker/host_header_reg_compute_checker.cpp+4-0
@@ -547,10 +547,14 @@ static void test_host_c_api_reg_compute_15()
547 using ::asc_loadalign_brc_postupdate_v3;547 using ::asc_loadalign_brc_postupdate_v3;
548 using ::asc_loadalign_brc_v2;548 using ::asc_loadalign_brc_v2;
549 using ::asc_loadalign_brc_v3;549 using ::asc_loadalign_brc_v3;
550+ using ::asc_loadalign_datablock_strided;
550 using ::asc_loadalign_deintlv;551 using ::asc_loadalign_deintlv;
551 using ::asc_loadalign_deintlv_postupdate;552 using ::asc_loadalign_deintlv_postupdate;
552 using ::asc_loadalign_downsample;553 using ::asc_loadalign_downsample;
553 using ::asc_loadalign_downsample_postupdate;554 using ::asc_loadalign_downsample_postupdate;
555+ using ::asc_loadalign_mask;
556+ using ::asc_loadalign_mask_downsample;
557+ using ::asc_loadalign_mask_upsample;
554 using ::asc_loadalign_postupdate;558 using ::asc_loadalign_postupdate;
555 using ::asc_loadalign_unpack;559 using ::asc_loadalign_unpack;
556 using ::asc_loadalign_unpack4;560 using ::asc_loadalign_unpack4;
Mtests/api/c_api/npu_arch_3510/vector_datamove/test_asc_load.cpp+49-3
@@ -8,6 +8,7 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11+#include <type_traits>
11#include <gtest/gtest.h>12#include <gtest/gtest.h>
12#include <mockcpp/mockcpp.hpp>13#include <mockcpp/mockcpp.hpp>
13#include "tests/api/c_api/stub/cce_stub.h"14#include "tests/api/c_api/stub/cce_stub.h"
@@ -23,15 +24,20 @@
23 }; \24 }; \
24 \25 \
25 namespace { \26 namespace { \
26- void cce_name1##_##data_type##_Stub1(vector_load_unalign& src0, __ubuf__ cce_data_type* src1) {} \27+ void cce_name1##_##data_type##_Stub1(vector_load_unalign& src0, __ubuf__ cce_data_type* src1) \
28+ { \
29+ EXPECT_EQ(src1, reinterpret_cast<__ubuf__ cce_data_type*>(11)); \
30+ } \
27 void cce_name0##_##data_type##_Stub0( \31 void cce_name0##_##data_type##_Stub0( \
28 vector_##cce_data_type& dst, vector_load_unalign& src0, __ubuf__ cce_data_type* src1) \32 vector_##cce_data_type& dst, vector_load_unalign& src0, __ubuf__ cce_data_type* src1) \
29- {} \33+ { \
34+ EXPECT_EQ(src1, reinterpret_cast<__ubuf__ cce_data_type*>(11)); \
35+ } \
30 } \36 } \
31 \37 \
32 TEST_F(TestVectorDataMove##class_name##_##data_type##_CApi, c_api_name##_##data_type##_Succ) \38 TEST_F(TestVectorDataMove##class_name##_##data_type##_CApi, c_api_name##_##data_type##_Succ) \
33 { \39 { \
34- __ubuf__ data_type* src = reinterpret_cast<__ubuf__ data_type*>(0); \40+ __ubuf__ data_type* src = reinterpret_cast<__ubuf__ data_type*>(11); \
35 dst_data_type dst; \41 dst_data_type dst; \
36 \42 \
37 MOCKER_CPP(cce_name1, void(vector_load_unalign&, __ubuf__ cce_data_type*)) \43 MOCKER_CPP(cce_name1, void(vector_load_unalign&, __ubuf__ cce_data_type*)) \
@@ -46,6 +52,26 @@
46 GlobalMockObject::verify(); \52 GlobalMockObject::verify(); \
47 }53 }
48 54 
55+#define TEST_VECTOR_DATAMOVE_LOAD_RETURN( \
56+ class_name, c_api_name, cce_name0, cce_name1, dst_data_type, data_type, cce_data_type) \
57+ TEST_F(TestVectorDataMove##class_name##_##data_type##_CApi, c_api_name##_##data_type##_ReturnSucc) \
58+ { \
59+ __ubuf__ data_type* src = reinterpret_cast<__ubuf__ data_type*>(11); \
60+ \
61+ MOCKER_CPP(cce_name1, void(vector_load_unalign&, __ubuf__ cce_data_type*)) \
62+ .times(1) \
63+ .will(invoke(cce_name1##_##data_type##_Stub1)); \
64+ \
65+ MOCKER_CPP(cce_name0, void(vector_##cce_data_type&, vector_store_unalign&, __ubuf__ cce_data_type*)) \
66+ .times(1) \
67+ .will(invoke(cce_name0##_##data_type##_Stub0)); \
68+ \
69+ static_assert(std::is_same_v<decltype(c_api_name(src)), dst_data_type>); \
70+ dst_data_type dst = c_api_name(src); \
71+ (void)dst; \
72+ GlobalMockObject::verify(); \
73+ }
74+ 
49TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_int8_t, int8_t, int8_t);75TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_int8_t, int8_t, int8_t);
50TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_uint8_t, uint8_t, uint8_t);76TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_uint8_t, uint8_t, uint8_t);
51TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_int16_t, int16_t, int16_t);77TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_int16_t, int16_t, int16_t);
@@ -63,3 +89,23 @@ TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_fp
63TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, fp4x2_e2m1_t);89TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, fp4x2_e2m1_t);
64TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, fp4x2_e1m2_t);90TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, fp4x2_e1m2_t);
65TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_int4x2_t, int4b_t, fp4x2_e1m2_t);91TEST_VECTOR_DATAMOVE_LOAD_INSTR(Vldasandvldus, asc_load, vldus, vldas, vector_int4x2_t, int4b_t, fp4x2_e1m2_t);
92+ 
93+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_int8_t, int8_t, int8_t);
94+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_uint8_t, uint8_t, uint8_t);
95+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_int16_t, int16_t, int16_t);
96+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_uint16_t, uint16_t, uint16_t);
97+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_int32_t, int32_t, int32_t);
98+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_uint32_t, uint32_t, uint32_t);
99+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_half, half, half);
100+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_float, float, float);
101+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_bfloat16_t, bfloat16_t, bfloat16_t);
102+TEST_VECTOR_DATAMOVE_LOAD_RETURN(
103+ Vldasandvldus, asc_load, vldus, vldas, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, fp8_e4m3fn_t);
104+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_hifloat8_t, hifloat8_t, uint8_t);
105+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_fp8_e5m2_t, fp8_e5m2_t, fp8_e5m2_t);
106+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_fp8_e8m0_t, fp8_e8m0_t, fp8_e8m0_t);
107+TEST_VECTOR_DATAMOVE_LOAD_RETURN(
108+ Vldasandvldus, asc_load, vldus, vldas, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, fp4x2_e2m1_t);
109+TEST_VECTOR_DATAMOVE_LOAD_RETURN(
110+ Vldasandvldus, asc_load, vldus, vldas, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, fp4x2_e1m2_t);
111+TEST_VECTOR_DATAMOVE_LOAD_RETURN(Vldasandvldus, asc_load, vldus, vldas, vector_int4x2_t, int4b_t, fp4x2_e1m2_t);
Mtests/api/c_api/npu_arch_3510/vector_datamove/test_asc_load_align_v2.cpp+72-24
@@ -8,35 +8,66 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11+#include <type_traits>
11#include <gtest/gtest.h>12#include <gtest/gtest.h>
12#include <mockcpp/mockcpp.hpp>13#include <mockcpp/mockcpp.hpp>
13#include "tests/api/c_api/stub/cce_stub.h"14#include "tests/api/c_api/stub/cce_stub.h"
14#include "include/c_api/asc_simd.h"15#include "include/c_api/asc_simd.h"
15 16 
16-#define TEST_VECTOR_COMPUTE_LOADALIGNV2(dst_type, src_type, cce_dst_type, cce_src_type) \17+namespace {
17- \18+constexpr uint16_t BLOCK_STRIDE = 3;
18- class TestVectorDataMoveLoadAlignV2##dst_type##CApi : public testing::Test { \19+constexpr uint16_t REPEAT_STRIDE = 5;
19- protected: \20+constexpr vector_bool LOAD_MASK = 0x5A5A;
20- void SetUp() {} \21+constexpr int32_t LOAD_OFFSET = (BLOCK_STRIDE << 16u) | REPEAT_STRIDE;
21- void TearDown() {} \22+} // namespace
22- }; \23+ 
23- \24+#define TEST_VECTOR_COMPUTE_LOADALIGNV2(dst_type, src_type, cce_dst_type, cce_src_type) \
24- namespace { \25+ \
25- void vsldb##_##dst_type##_Stub(cce_dst_type& dst, __ubuf__ cce_src_type* src, int32_t offset, vector_bool mask) {} \26+ class TestVectorDataMoveLoadAlignV2##dst_type##CApi : public testing::Test { \
26- } \27+ protected: \
27- \28+ void SetUp() {} \
28- TEST_F(TestVectorDataMoveLoadAlignV2##dst_type##CApi, c_api_asc_loadalign_v2_##dst_type##_Succ) \29+ void TearDown() {} \
29- { \30+ }; \
30- dst_type dst; \31+ \
31- __ubuf__ src_type* src = reinterpret_cast<__ubuf__ src_type*>(11); \32+ namespace { \
32- vector_bool mask; \33+ void vsldb##_##dst_type##_Stub(cce_dst_type& dst, __ubuf__ cce_src_type* src, int32_t offset, vector_bool mask) \
33- \34+ { \
34- MOCKER_CPP(vsldb, void(cce_dst_type&, __ubuf__ cce_src_type*, int32_t, vector_bool)) \35+ EXPECT_EQ(src, reinterpret_cast<__ubuf__ cce_src_type*>(11)); \
35- .times(1) \36+ EXPECT_EQ(offset, LOAD_OFFSET); \
36- .will(invoke(vsldb##_##dst_type##_Stub)); \37+ EXPECT_EQ(mask, LOAD_MASK); \
37- \38+ } \
38- asc_loadalign(dst, src, 0, 0, mask); \39+ } \
39- GlobalMockObject::verify(); \40+ \
41+ TEST_F(TestVectorDataMoveLoadAlignV2##dst_type##CApi, c_api_asc_loadalign_v2_##dst_type##_Succ) \
42+ { \
43+ dst_type dst; \
44+ __ubuf__ src_type* src = reinterpret_cast<__ubuf__ src_type*>(11); \
45+ vector_bool mask = LOAD_MASK; \
46+ \
47+ MOCKER_CPP(vsldb, void(cce_dst_type&, __ubuf__ cce_src_type*, int32_t, vector_bool)) \
48+ .times(1) \
49+ .will(invoke(vsldb##_##dst_type##_Stub)); \
50+ \
51+ asc_loadalign(dst, src, BLOCK_STRIDE, REPEAT_STRIDE, mask); \
52+ GlobalMockObject::verify(); \
53+ }
54+ 
55+#define TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(dst_type, src_type, cce_dst_type, cce_src_type) \
56+ TEST_F( \
57+ TestVectorDataMoveLoadAlignV2##dst_type##CApi, c_api_asc_loadalign_datablock_strided_##dst_type##_ReturnSucc) \
58+ { \
59+ __ubuf__ src_type* src = reinterpret_cast<__ubuf__ src_type*>(11); \
60+ vector_bool mask = LOAD_MASK; \
61+ \
62+ MOCKER_CPP(vsldb, void(cce_dst_type&, __ubuf__ cce_src_type*, int32_t, vector_bool)) \
63+ .times(1) \
64+ .will(invoke(vsldb##_##dst_type##_Stub)); \
65+ \
66+ static_assert(std::is_same_v< \
67+ decltype(asc_loadalign_datablock_strided(src, BLOCK_STRIDE, REPEAT_STRIDE, mask)), dst_type>); \
68+ dst_type dst = asc_loadalign_datablock_strided(src, BLOCK_STRIDE, REPEAT_STRIDE, mask); \
69+ (void)dst; \
70+ GlobalMockObject::verify(); \
40 }71 }
41 72 
42TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_int8_t, int8_t, vector_int8_t, int8_t);73TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_int8_t, int8_t, vector_int8_t, int8_t);
@@ -57,6 +88,23 @@ TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_half, half, vector_half, half);
57TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_float, float, vector_float, float);88TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_float, float, vector_float, float);
58TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t, float4_e1m2x2_t);89TEST_VECTOR_COMPUTE_LOADALIGNV2(vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t, float4_e1m2x2_t);
59 90 
91+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_int8_t, int8_t, vector_int8_t, int8_t);
92+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
93+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_int16_t, int16_t, vector_int16_t, int16_t);
94+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_uint16_t, uint16_t, vector_uint16_t, uint16_t);
95+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_int32_t, int32_t, vector_int32_t, int32_t);
96+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_uint32_t, uint32_t, vector_uint32_t, uint32_t);
97+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_bfloat16_t, bfloat16_t, vector_bfloat16_t, bfloat16_t);
98+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t, fp8_e4m3fn_t);
99+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
100+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t, fp8_e5m2_t);
101+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t, fp8_e8m0_t);
102+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t, fp4x2_e1m2_t);
103+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t, fp4x2_e2m1_t);
104+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_half, half, vector_half, half);
105+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_float, float, vector_float, float);
106+TEST_VECTOR_COMPUTE_LOADALIGN_RETURN(vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t, float4_e1m2x2_t);
107+ 
60#define TEST_VECTOR_COMPUTE_LOADALIGN_POSTUPDATEV2(dst_type, src_type, cce_dst_type, cce_src_type) \108#define TEST_VECTOR_COMPUTE_LOADALIGN_POSTUPDATEV2(dst_type, src_type, cce_dst_type, cce_src_type) \
61 \109 \
62 class TestVectorDataMoveLoadAlignPostUpdateV2##dst_type##CApi : public testing::Test { \110 class TestVectorDataMoveLoadAlignPostUpdateV2##dst_type##CApi : public testing::Test { \
Atests/api/c_api/npu_arch_3510/vector_datamove/test_asc_load_return.cpp+396-0
@@ -0,0 +1,396 @@
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 <type_traits>
12+#include <gtest/gtest.h>
13+#include <mockcpp/mockcpp.hpp>
14+#include "tests/api/c_api/stub/cce_stub.h"
15+#include "include/c_api/asc_simd.h"
16+ 
17+namespace {
18+template <typename T>
19+constexpr Literal LoadAlignReturnExpectedDist()
20+{
21+ return Literal::NORM;
22+}
23+ 
24+template <typename T>
25+constexpr Literal LoadAlignBrcElemReturnExpectedDist()
26+{
27+ if constexpr (sizeof(T) == 1) {
28+ return Literal::BRC_B8;
29+ } else if constexpr (sizeof(T) == 2) {
30+ return Literal::BRC_B16;
31+ } else {
32+ return Literal::BRC_B32;
33+ }
34+}
35+ 
36+template <typename T>
37+constexpr Literal LoadAlignBrcDatablockReturnExpectedDist()
38+{
39+ return Literal::BLK;
40+}
41+ 
42+template <typename T>
43+constexpr Literal LoadAlignBrcElem2DatablockReturnExpectedDist()
44+{
45+ if constexpr (sizeof(T) == 2) {
46+ return Literal::E2B_B16;
47+ } else {
48+ return Literal::E2B_B32;
49+ }
50+}
51+ 
52+template <typename T>
53+constexpr Literal LoadAlignDownsampleReturnExpectedDist()
54+{
55+ if constexpr (sizeof(T) == 1) {
56+ return Literal::DS_B8;
57+ } else {
58+ return Literal::DS_B16;
59+ }
60+}
61+ 
62+template <typename T>
63+constexpr Literal LoadAlignMaskDownsampleReturnExpectedDist()
64+{
65+ return Literal::DS;
66+}
67+ 
68+template <typename T>
69+constexpr Literal LoadAlignUnpackReturnExpectedDist()
70+{
71+ if constexpr (sizeof(T) == 1) {
72+ return Literal::UNPK_B8;
73+ } else if constexpr (sizeof(T) == 2) {
74+ return Literal::UNPK_B16;
75+ } else {
76+ return Literal::UNPK_B32;
77+ }
78+}
79+ 
80+template <typename T>
81+constexpr Literal LoadAlignUnpack4ReturnExpectedDist()
82+{
83+ return Literal::UNPK4_B8;
84+}
85+ 
86+template <typename T>
87+constexpr Literal LoadAlignUpsampleReturnExpectedDist()
88+{
89+ if constexpr (sizeof(T) == 1) {
90+ return Literal::US_B8;
91+ } else {
92+ return Literal::US_B16;
93+ }
94+}
95+ 
96+template <typename T>
97+constexpr Literal LoadAlignMaskUpsampleReturnExpectedDist()
98+{
99+ return Literal::US;
100+}
101+} // namespace
102+ 
103+#define TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN( \
104+ class_name, c_api_name, cce_name, dst_type, src_type, cce_dst_type, cce_src_type) \
105+ class TestVectorDatamove##class_name##dst_type##src_type##CApi : public testing::Test { \
106+ protected: \
107+ void SetUp() {} \
108+ void TearDown() {} \
109+ }; \
110+ \
111+ namespace { \
112+ void c_api_name##_##dst_type##_##src_type##_ReturnStub( \
113+ cce_dst_type& dst, __ubuf__ cce_src_type* src, int32_t offset, Literal load_dist) \
114+ { \
115+ EXPECT_EQ(src, reinterpret_cast<__ubuf__ cce_src_type*>(11)); \
116+ EXPECT_EQ(offset, static_cast<int32_t>(0)); \
117+ EXPECT_EQ(load_dist, class_name##ExpectedDist<src_type>()); \
118+ } \
119+ } \
120+ \
121+ TEST_F( \
122+ TestVectorDatamove##class_name##dst_type##src_type##CApi, c_api_name##_##dst_type##_##src_type##_ReturnSucc) \
123+ { \
124+ __ubuf__ src_type* src = reinterpret_cast<__ubuf__ src_type*>(11); \
125+ \
126+ MOCKER_CPP(cce_name, void(cce_dst_type&, __ubuf__ cce_src_type*, int32_t, Literal)) \
127+ .times(1) \
128+ .will(invoke(c_api_name##_##dst_type##_##src_type##_ReturnStub)); \
129+ \
130+ static_assert(std::is_same_v<decltype(c_api_name(src)), dst_type>); \
131+ dst_type dst = c_api_name(src); \
132+ (void)dst; \
133+ GlobalMockObject::verify(); \
134+ }
135+ 
136+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
137+ LoadAlignReturn, asc_loadalign, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
138+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
139+ LoadAlignReturn, asc_loadalign, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
140+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
141+ LoadAlignReturn, asc_loadalign, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t, fp4x2_e2m1_t);
142+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
143+ LoadAlignReturn, asc_loadalign, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t, fp4x2_e1m2_t);
144+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
145+ LoadAlignReturn, asc_loadalign, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t, float4_e1m2x2_t);
146+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
147+ LoadAlignReturn, asc_loadalign, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t, fp8_e8m0_t);
148+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
149+ LoadAlignReturn, asc_loadalign, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t, fp8_e5m2_t);
150+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
151+ LoadAlignReturn, asc_loadalign, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t, fp8_e4m3fn_t);
152+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
153+ LoadAlignReturn, asc_loadalign, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
154+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
155+ LoadAlignReturn, asc_loadalign, vlds, vector_int16_t, int16_t, vector_int16_t, int16_t);
156+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
157+ LoadAlignReturn, asc_loadalign, vlds, vector_uint16_t, uint16_t, vector_uint16_t, uint16_t);
158+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(LoadAlignReturn, asc_loadalign, vlds, vector_half, half, vector_half, half);
159+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
160+ LoadAlignReturn, asc_loadalign, vlds, vector_bfloat16_t, bfloat16_t, vector_bfloat16_t, bfloat16_t);
161+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
162+ LoadAlignReturn, asc_loadalign, vlds, vector_int32_t, int32_t, vector_int32_t, int32_t);
163+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
164+ LoadAlignReturn, asc_loadalign, vlds, vector_uint32_t, uint32_t, vector_uint32_t, uint32_t);
165+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(LoadAlignReturn, asc_loadalign, vlds, vector_float, float, vector_float, float);
166+ 
167+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
168+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
169+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
170+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
171+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
172+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t,
173+ fp4x2_e2m1_t);
174+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
175+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t,
176+ fp4x2_e1m2_t);
177+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
178+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t,
179+ float4_e1m2x2_t);
180+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
181+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t, fp8_e8m0_t);
182+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
183+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t, fp8_e5m2_t);
184+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
185+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t,
186+ fp8_e4m3fn_t);
187+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
188+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
189+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
190+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_int16_t, int16_t, vector_int16_t, int16_t);
191+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
192+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_uint16_t, uint16_t, vector_uint16_t, uint16_t);
193+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
194+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_half, half, vector_half, half);
195+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
196+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_bfloat16_t, bfloat16_t, vector_bfloat16_t, bfloat16_t);
197+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
198+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_int32_t, int32_t, vector_int32_t, int32_t);
199+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
200+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_uint32_t, uint32_t, vector_uint32_t, uint32_t);
201+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
202+ LoadAlignBrcElemReturn, asc_loadalign_brc_elem, vlds, vector_float, float, vector_float, float);
203+ 
204+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
205+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
206+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
207+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
208+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
209+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t,
210+ vector_fp4x2_e2m1_t, fp4x2_e2m1_t);
211+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
212+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t,
213+ vector_fp4x2_e1m2_t, fp4x2_e1m2_t);
214+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
215+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t,
216+ float4_e1m2x2_t);
217+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
218+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t,
219+ fp8_e8m0_t);
220+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
221+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t,
222+ fp8_e5m2_t);
223+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
224+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t,
225+ vector_fp8_e4m3fn_t, fp8_e4m3fn_t);
226+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
227+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t,
228+ uint8_t);
229+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
230+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_int16_t, int16_t, vector_int16_t, int16_t);
231+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
232+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_uint16_t, uint16_t, vector_uint16_t,
233+ uint16_t);
234+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
235+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_half, half, vector_half, half);
236+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
237+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_bfloat16_t, bfloat16_t, vector_bfloat16_t,
238+ bfloat16_t);
239+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
240+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_int32_t, int32_t, vector_int32_t, int32_t);
241+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
242+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_uint32_t, uint32_t, vector_uint32_t,
243+ uint32_t);
244+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
245+ LoadAlignBrcDatablockReturn, asc_loadalign_brc_datablock, vlds, vector_float, float, vector_float, float);
246+ 
247+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
248+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_int16_t, int16_t, vector_int16_t,
249+ int16_t);
250+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
251+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_uint16_t, uint16_t,
252+ vector_uint16_t, uint16_t);
253+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
254+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_half, half, vector_half, half);
255+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
256+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_bfloat16_t, bfloat16_t,
257+ vector_bfloat16_t, bfloat16_t);
258+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
259+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_int32_t, int32_t, vector_int32_t,
260+ int32_t);
261+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
262+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_uint32_t, uint32_t,
263+ vector_uint32_t, uint32_t);
264+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
265+ LoadAlignBrcElem2DatablockReturn, asc_loadalign_brc_elem2datablock, vlds, vector_float, float, vector_float, float);
266+ 
267+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
268+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
269+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
270+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
271+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
272+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t,
273+ fp4x2_e2m1_t);
274+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
275+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t,
276+ fp4x2_e1m2_t);
277+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
278+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t,
279+ float4_e1m2x2_t);
280+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
281+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t,
282+ fp8_e8m0_t);
283+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
284+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t,
285+ fp8_e5m2_t);
286+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
287+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t,
288+ fp8_e4m3fn_t);
289+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
290+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
291+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
292+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_int16_t, int16_t, vector_int16_t, int16_t);
293+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
294+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_uint16_t, uint16_t, vector_uint16_t, uint16_t);
295+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
296+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_half, half, vector_half, half);
297+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
298+ LoadAlignDownsampleReturn, asc_loadalign_downsample, vlds, vector_bfloat16_t, bfloat16_t, vector_bfloat16_t,
299+ bfloat16_t);
300+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
301+ LoadAlignMaskDownsampleReturn, asc_loadalign_mask_downsample, plds, vector_bool, uint32_t, vector_bool, uint32_t);
302+ 
303+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
304+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
305+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
306+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
307+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
308+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t,
309+ fp4x2_e2m1_t);
310+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
311+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t,
312+ fp4x2_e1m2_t);
313+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
314+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t, float4_e1m2x2_t);
315+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
316+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t, fp8_e8m0_t);
317+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
318+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t, fp8_e5m2_t);
319+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
320+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t,
321+ fp8_e4m3fn_t);
322+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
323+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
324+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
325+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_int16_t, int16_t, vector_int16_t, int16_t);
326+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
327+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_uint16_t, uint16_t, vector_uint16_t, uint16_t);
328+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
329+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_half, half, vector_half, half);
330+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
331+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_bfloat16_t, bfloat16_t, vector_bfloat16_t, bfloat16_t);
332+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
333+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_int32_t, int32_t, vector_int32_t, int32_t);
334+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
335+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_uint32_t, uint32_t, vector_uint32_t, uint32_t);
336+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
337+ LoadAlignUnpackReturn, asc_loadalign_unpack, vlds, vector_float, float, vector_float, float);
338+ 
339+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
340+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
341+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
342+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
343+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
344+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t,
345+ fp4x2_e2m1_t);
346+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
347+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t,
348+ fp4x2_e1m2_t);
349+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
350+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t,
351+ float4_e1m2x2_t);
352+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
353+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t, fp8_e8m0_t);
354+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
355+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t, fp8_e5m2_t);
356+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
357+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t,
358+ fp8_e4m3fn_t);
359+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
360+ LoadAlignUnpack4Return, asc_loadalign_unpack4, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
361+ 
362+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
363+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_int8_t, int8_t, vector_int8_t, int8_t);
364+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
365+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_uint8_t, uint8_t, vector_uint8_t, uint8_t);
366+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
367+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_fp4x2_e2m1_t, fp4x2_e2m1_t, vector_fp4x2_e2m1_t,
368+ fp4x2_e2m1_t);
369+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
370+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_fp4x2_e1m2_t, fp4x2_e1m2_t, vector_fp4x2_e1m2_t,
371+ fp4x2_e1m2_t);
372+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
373+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_int4x2_t, int4b_t, vector_fp4x2_e1m2_t,
374+ float4_e1m2x2_t);
375+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
376+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_fp8_e8m0_t, fp8_e8m0_t, vector_fp8_e8m0_t,
377+ fp8_e8m0_t);
378+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
379+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_fp8_e5m2_t, fp8_e5m2_t, vector_fp8_e5m2_t,
380+ fp8_e5m2_t);
381+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
382+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_fp8_e4m3fn_t, fp8_e4m3fn_t, vector_fp8_e4m3fn_t,
383+ fp8_e4m3fn_t);
384+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
385+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_hifloat8_t, hifloat8_t, vector_uint8_t, uint8_t);
386+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
387+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_int16_t, int16_t, vector_int16_t, int16_t);
388+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
389+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_uint16_t, uint16_t, vector_uint16_t, uint16_t);
390+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
391+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_half, half, vector_half, half);
392+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
393+ LoadAlignUpsampleReturn, asc_loadalign_upsample, vlds, vector_bfloat16_t, bfloat16_t, vector_bfloat16_t,
394+ bfloat16_t);
395+TEST_VECTOR_DATAMOVE_LOADALIGN_RETURN(
396+ LoadAlignMaskUpsampleReturn, asc_loadalign_mask_upsample, plds, vector_bool, uint32_t, vector_bool, uint32_t);
Mtests/api/c_api/npu_arch_3510/vector_datamove/test_asc_loadalign_plds.cpp+13-0
@@ -8,6 +8,7 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11+#include <type_traits>
11#include <gtest/gtest.h>12#include <gtest/gtest.h>
12#include <mockcpp/mockcpp.hpp>13#include <mockcpp/mockcpp.hpp>
13#include "tests/api/c_api/stub/cce_stub.h"14#include "tests/api/c_api/stub/cce_stub.h"
@@ -100,6 +101,18 @@ TEST_F(TestVectorDatamoveLoadAlignPlds, LoadAlignPlds_Succ)
100 GlobalMockObject::verify();101 GlobalMockObject::verify();
101}102}
102 103 
104+TEST_F(TestVectorDatamoveLoadAlignPlds, LoadAlignMaskPlds_ReturnSucc)
105+{
106+ __ubuf__ uint32_t* src = reinterpret_cast<__ubuf__ uint32_t*>(22);
107+ 
108+ MOCKER_CPP(plds, void(vector_bool&, __ubuf__ uint32_t*, int32_t, Literal)).times(1).will(invoke(plds_stub));
109+ 
110+ static_assert(std::is_same_v<decltype(asc_loadalign_mask(src)), vector_bool>);
111+ vector_bool dst = asc_loadalign_mask(src);
112+ (void)dst;
113+ GlobalMockObject::verify();
114+}
115+ 
103TEST_F(TestVectorDatamoveLoadAlignPlds, LoadAlignPldsUpsample_Succ)116TEST_F(TestVectorDatamoveLoadAlignPlds, LoadAlignPldsUpsample_Succ)
104{117{
105 vector_bool dst = asc_create_mask_b16(PAT_ALL);118 vector_bool dst = asc_create_mask_b16(PAT_ALL);