docs: 完善C API数据依赖同步约束说明-asc-devkit-AtomGit

Merged

docs: 完善C API数据依赖同步约束说明 #6689

lihuaichaoAdd on 26 days agoMerged fromlihuaichao/asc-devkit_950： docs_capi_constraint into cann/asc-devkit： master

docs: 完善C API数据依赖同步约束说明 #6689

Merged

lihuaichaoAdd on 26 days ago

Merged fromlihuaichao/asc-devkit_950 docs_capi_constraintintocann/asc-devkit master

Collapse

docs / zh / api / SIMD-API / c_api
- cube_datamove
  - cube_compute_load
    - asc_copy_gm2l1_2d_arch_3510.md
    - asc_copy_gm2l1_align.md
    - asc_copy_gm2l1_arch_2201.md
    - asc_copy_gm2l1_dn2nz.md
    - asc_copy_gm2l1_highdim_split_arch_3510.md
    - asc_copy_gm2l1_nd2nz_arch_2201.md
    - asc_copy_gm2l1_nd2nz_arch_3510.md
    - asc_copy_l12bt_arch_3510.md
    - asc_copy_l12fb_arch_3510.md
    - asc_copy_l12l0a_2d_arch_3510.md
    - asc_copy_l12l0a_3d_arch_3510.md
    - asc_copy_l12l0a_mx.md
    - asc_copy_l12l0b_2d_arch_3510.md
    - asc_copy_l12l0b_3d_arch_3510.md
    - asc_copy_l12l0b_mx.md
    - asc_copy_l12l0b_trans_arch_3510.md
    - asc_copy_l12l0c.md
  - cube_compute_store
    - asc_copy_l0c2gm_arch_3510.md
    - asc_copy_l0c2l1_arch_2201.md
    - asc_copy_l0c2l1_arch_3510.md
    - asc_copy_l0c2ub.md
    - asc_copy_l12gm.md
    - asc_copy_l12ub.md
  - cube_load_aux_config
    - asc_fill_l1.md
- reg_compute
  - load
    - asc_load.md
    - asc_loadalign.md
    - asc_loadalign_brc_datablock.md
    - asc_loadalign_brc_datablock_postupdate.md
    - asc_loadalign_brc_elem.md
    - asc_loadalign_brc_elem2datablock.md
    - asc_loadalign_brc_elem2datablock_postupdate.md
    - asc_loadalign_brc_elem_postupdate.md
    - asc_loadalign_datablock_strided.md
    - asc_loadalign_deintlv.md
    - asc_loadalign_deintlv_postupdate.md
    - asc_loadalign_downsample.md
    - asc_loadalign_downsample_postupdate.md
    - asc_loadalign_mask.md
    - asc_loadalign_mask_downsample.md
    - asc_loadalign_mask_upsample.md
    - asc_loadalign_postupdate.md
    - asc_loadalign_unpack.md
    - asc_loadalign_unpack4.md
    - asc_loadalign_unpack4_postupdate.md
    - asc_loadalign_unpack_postupdate.md
    - asc_loadalign_upsample.md
    - asc_loadalign_upsample_postupdate.md
    - asc_loadunalign.md
    - asc_loadunalign_postupdate.md
    - asc_loadunalign_pre.md
    - reg_load_overview.md
  - reg_mask
    - asc_get_mask_spr.md
  - reg_permute_sel
    - asc_squeeze_and_storeunalign.md
    - asc_squeeze_and_storeunalign_finalize.md
  - scatter
    - asc_scatter.md
  - store
    - asc_store.md
    - asc_storealign.md
    - asc_storealign_1st.md
    - asc_storealign_1st_postupdate.md
    - asc_storealign_intlv.md
    - asc_storealign_pack.md
    - asc_storealign_pack_postupdate.md
    - asc_storealign_pack_quarter.md
    - asc_storealign_pack_quarter_postupdate.md
    - asc_storealign_postupdate.md
    - asc_storealign_squeeze_status.md
    - asc_storealign_squeeze_status_postupdate.md
    - asc_storeunalign.md
    - asc_storeunalign_post.md
    - asc_storeunalign_post_postupdate.md
    - asc_storeunalign_postupdate.md
- vector_datamove
  - asc_copy_gm2ub
    - asc_copy_gm2ub_arch_2201.md
    - asc_copy_gm2ub_arch_3510.md
  - asc_copy_gm2ub_align
    - asc_copy_gm2ub_align_arch_2201.md
    - asc_copy_gm2ub_align_arch_3510.md
  - asc_copy_ub2gm
    - asc_copy_ub2gm_arch_2201.md
    - asc_copy_ub2gm_arch_3510.md
  - asc_copy_ub2gm_align
    - asc_copy_ub2gm_align_arch_2201.md
    - asc_copy_ub2gm_align_arch_3510.md
  - asc_copy_ub2l1.md
  - asc_copy_ub2ub.md
  - asc_ndim_copy_gm2ub.md

Total: 82 files changed+122-119

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_2d_arch_3510.md+1-1

	@@ -96,7 +96,7 @@ PIPE_MTE2
96	- 本接口非AIC调用直接返回。	96	- 本接口非AIC调用直接返回。
97	- dst起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	97	- dst起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
98	- src起始地址需1字节对齐（Global Memory对齐要求），否则会导致搬运异常。	98	- src起始地址需1字节对齐（Global Memory对齐要求），否则会导致搬运异常。
99	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	99	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
		L L Lycheeeee26 days ago 修正搬运的非对齐相关描述。
100	- L1 Buffer容量上限：L1 Buffer总容量512KB，dst偏移量与搬运大小之和不可越界，否则触发异常。	100	- L1 Buffer容量上限：L1 Buffer总容量512KB，dst偏移量与搬运大小之和不可越界，否则触发异常。
101		101
102	### 2D矩阵搬运模式约束	102	### 2D矩阵搬运模式约束

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_align.md+2-2

	@@ -28,7 +28,7 @@
28		28
29	头文件路径为：`"c_api/cube_datamove/cube_datamove.h"`。	29	头文件路径为：`"c_api/cube_datamove/cube_datamove.h"`。
30		30
31	-将数据从Global Memory搬运到L1 Buffer，提供非32B对齐搬运能力，搬运过程中对支持对数据进行padding。	31	+将数据从Global Memory（GM）搬运到L1 Buffer，支持搬运长度不是32字节整数倍的数据，并在搬运过程中进行32字节对齐补齐。搬运数据长度仍需满足对应数据类型的字节对齐要求；源、目的起始地址的对齐要求请参见[参数说明](#参数说明)。
32		32
33	若搬运数据长度非32字节对齐，会将搬运数据补齐至32字节对齐。可通过配置参数`dst_stride`选择Normal模式或Compact模式。非32字节对齐场景支持以下两种填充方式：	33	若搬运数据长度非32字节对齐，会将搬运数据补齐至32字节对齐。可通过配置参数`dst_stride`选择Normal模式或Compact模式。非32字节对齐场景支持以下两种填充方式：
34		34
	@@ -132,7 +132,7 @@ PIPE_MTE2
132	- 本接口非AIC调用直接返回。	132	- 本接口非AIC调用直接返回。
133	- `dst`起始地址需32字节对齐（L1 Buffer对齐要求），否则触发异常。	133	- `dst`起始地址需32字节对齐（L1 Buffer对齐要求），否则触发异常。
134	- `src`起始地址需1字节对齐（Global Memory对齐要求），不满足触发异常。	134	- `src`起始地址需1字节对齐（Global Memory对齐要求），不满足触发异常。
135	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	135	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
136	- L1 Buffer容量上限：L1 Buffer总容量512KB，dst偏移与搬运大小之和不可越界，否则触发目的地址越界异常。	136	- L1 Buffer容量上限：L1 Buffer总容量512KB，dst偏移与搬运大小之和不可越界，否则触发目的地址越界异常。
137	- `burst_len`、`burst_count`、`src_stride`、`dst_stride`、`left_padding_count`与`right_padding_count`需满足参数说明的取值范围，否则会导致搬运结果不符合预期。	137	- `burst_len`、`burst_count`、`src_stride`、`dst_stride`、`left_padding_count`与`right_padding_count`需满足参数说明的取值范围，否则会导致搬运结果不符合预期。
138	- `enable_data_select`设置为true时，或者左右填充模式（`left_padding_count`和`right_padding_count`任意不为0）时，须先调用[asc_set_gm2l1_pad](../cube_load_aux_config/asc_set_gm2l1_pad.md)配置填充值。	138	- `enable_data_select`设置为true时，或者左右填充模式（`left_padding_count`和`right_padding_count`任意不为0）时，须先调用[asc_set_gm2l1_pad](../cube_load_aux_config/asc_set_gm2l1_pad.md)配置填充值。

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_arch_2201.md+1-1

	@@ -148,7 +148,7 @@ PIPE_MTE2
148	## 约束说明	148	## 约束说明
149		149
150	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	150	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
151	-- 如果需要执行多条asc_copy_gm2l1指令，且asc_copy_gm2l1指令的目的地址存在重叠，需要插入同步指令，保证多个asc_copy_gm2l1指令的串行化，防止出现异常数据。	151	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
152		152
153	## 调用示例	153	## 调用示例
154		154

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_dn2nz.md+1-1

	@@ -104,7 +104,7 @@ PIPE_MTE2
104	- 本接口非AIC调用直接返回。	104	- 本接口非AIC调用直接返回。
105	- `dst`的起始地址要求32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	105	- `dst`的起始地址要求32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
106	- `src`的起始地址要求1字节对齐（Global Memory对齐要求），否则会导致搬运异常。	106	- `src`的起始地址要求1字节对齐（Global Memory对齐要求），否则会导致搬运异常。
107	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	107	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
108	- L1 Buffer容量上限：L1 Buffer总容量512KB，`dst`偏移量与搬运大小之和不可越界，否则触发异常。	108	- L1 Buffer容量上限：L1 Buffer总容量512KB，`dst`偏移量与搬运大小之和不可越界，否则触发异常。
109		109
110	### DN2Nz约束	110	### DN2Nz约束

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_highdim_split_arch_3510.md+1-1

	@@ -83,7 +83,7 @@ PIPE_MTE2
83	- 本接口仅在AIC上生效，非AIC调用直接返回。	83	- 本接口仅在AIC上生效，非AIC调用直接返回。
84	- `dst`起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	84	- `dst`起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
85	- `src`起始地址需1字节对齐（Global Memory对齐要求），否则会导致搬运异常。	85	- `src`起始地址需1字节对齐（Global Memory对齐要求），否则会导致搬运异常。
86	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	86	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
87	- L1 Buffer容量上限：L1 Buffer总容量512KB，dst偏移量与搬运大小之和不可越界，否则触发异常。	87	- L1 Buffer容量上限：L1 Buffer总容量512KB，dst偏移量与搬运大小之和不可越界，否则触发异常。
88		88
89	### 高维切分数据搬运约束	89	### 高维切分数据搬运约束

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_nd2nz_arch_2201.md+1-1

	@@ -91,7 +91,7 @@ PIPE_MTE2
91	## 约束说明	91	## 约束说明
92		92
93	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	93	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
94	-- 如果需要执行多条asc_copy_gm2l1_nd2nz指令，且asc_copy_gm2l1_nd2nz指令的目的地址存在重叠，需要插入同步指令，保证多个asc_copy_gm2l1_nd2nz指令的串行化，防止出现异常数据。	94	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
95		95
96	## 调用示例	96	## 调用示例
97		97

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_gm2l1_nd2nz_arch_3510.md+1-1

	@@ -109,7 +109,7 @@ PIPE_MTE2
109	- 本接口非AIC调用直接返回。	109	- 本接口非AIC调用直接返回。
110	- `dst`起始地址需要按照32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	110	- `dst`起始地址需要按照32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
111	- `src`起始地址需要按照1字节对齐（GM对齐要求），否则会导致搬运异常。	111	- `src`起始地址需要按照1字节对齐（GM对齐要求），否则会导致搬运异常。
112	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	112	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
113	- L1 Buffer容量上限：L1 Buffer总容量512KB，`dst`偏移量与搬运大小之和不可越界，否则触发异常。	113	- L1 Buffer容量上限：L1 Buffer总容量512KB，`dst`偏移量与搬运大小之和不可越界，否则触发异常。
114		114
115	### ND2Nz搬运约束	115	### ND2Nz搬运约束

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12bt_arch_3510.md+1-1

	@@ -124,7 +124,7 @@ PIPE_MTE1
124	- dst起始地址需64字节对齐（BiasTable Buffer对齐要求），src起始地址需32字节对齐（L1 Buffer对齐要求），否则触发异常。	124	- dst起始地址需64字节对齐（BiasTable Buffer对齐要求），src起始地址需32字节对齐（L1 Buffer对齐要求），否则触发异常。
125	- BiasTable Buffer容量上限4KB，dst偏移与搬运大小之和不可越界，否则触发写溢出异常。	125	- BiasTable Buffer容量上限4KB，dst偏移与搬运大小之和不可越界，否则触发写溢出异常。
126	- L1 Buffer容量上限512KB，src偏移与搬运大小之和不可越界，否则触发读异常。	126	- L1 Buffer容量上限512KB，src偏移与搬运大小之和不可越界，否则触发读异常。
127	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	127	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
128	- n_burst为0时不执行搬运，本接口被视为NOP（空操作）。	128	- n_burst为0时不执行搬运，本接口被视为NOP（空操作）。
129	- dtype为`float`或`int32_t`时，len_burst与dst_gap均必须2对齐（即取值为偶数），size/32需为偶数，否则触发搬运异常。	129	- dtype为`float`或`int32_t`时，len_burst与dst_gap均必须2对齐（即取值为偶数），size/32需为偶数，否则触发搬运异常。
130	- conv_control开启（非0）时dtype必须为`half`，否则触发搬运异常，传入非0/1的值仅最低位生效（如2等价于0即转换关闭、3等价于1即转换开启）。	130	- conv_control开启（非0）时dtype必须为`half`，否则触发搬运异常，传入非0/1的值仅最低位生效（如2等价于0即转换关闭、3等价于1即转换开启）。

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12fb_arch_3510.md+1-1

	@@ -89,7 +89,7 @@ PIPE_FIX
89	- dst起始地址根据功能不同对齐要求不同，对于随路量化场景，起始地址需要128B对齐，对于随路Relu场景起始地址需要64B对齐，否则触发搬运异常。	89	- dst起始地址根据功能不同对齐要求不同，对于随路量化场景，起始地址需要128B对齐，对于随路Relu场景起始地址需要64B对齐，否则触发搬运异常。
90	- L1 Buffer总容量512KB，src偏移与搬运大小之和不可越界，否则触发地址越界异常。	90	- L1 Buffer总容量512KB，src偏移与搬运大小之和不可越界，否则触发地址越界异常。
91	- Fixpipe Buffer对于随路量化场景，容量上限为4KB；对于随路Relu场景，容量上限为2KB。dst偏移与搬运大小之和越界时会截断处理，导致搬运结果不符合预期。	91	- Fixpipe Buffer对于随路量化场景，容量上限为4KB；对于随路Relu场景，容量上限为2KB。dst偏移与搬运大小之和越界时会截断处理，导致搬运结果不符合预期。
92	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	92	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
93	- 对于随路量化场景，`len_burst`与`dst_gap_size`场景取值必须为2的倍数，否则触发搬运异常。	93	- 对于随路量化场景，`len_burst`与`dst_gap_size`场景取值必须为2的倍数，否则触发搬运异常。
94	- `n_burst`或`size`设置为`0`时，不执行搬运，本接口被视为NOP（空操作）。	94	- `n_burst`或`size`设置为`0`时，不执行搬运，本接口被视为NOP（空操作）。
95		95

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0a_2d_arch_3510.md+1-1

	@@ -149,7 +149,7 @@ PIPE_MTE1
149	- dst起始地址需要按照512字节对齐（L0A Buffer对齐要求），否则会导致搬运异常。	149	- dst起始地址需要按照512字节对齐（L0A Buffer对齐要求），否则会导致搬运异常。
150	- src起始地址需要按照32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	150	- src起始地址需要按照32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
151	- src位于L1 Buffer，dst位于L0A Buffer，两者分属不同的物理存储单元，不存在源操作数与目的操作数地址重叠的场景。	151	- src位于L1 Buffer，dst位于L0A Buffer，两者分属不同的物理存储单元，不存在源操作数与目的操作数地址重叠的场景。
152	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	152	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
153	- L0A Buffer容量上限：L0A Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发异常。	153	- L0A Buffer容量上限：L0A Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发异常。
154	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。	154	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。
155		155

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0a_3d_arch_3510.md+1-1

	@@ -132,7 +132,7 @@ PIPE_MTE1
132	- 本接口非AIC调用直接返回。	132	- 本接口非AIC调用直接返回。
133	- dst起始地址需512字节对齐（L0A Buffer对齐要求），否则会导致搬运异常。	133	- dst起始地址需512字节对齐（L0A Buffer对齐要求），否则会导致搬运异常。
134	- src起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	134	- src起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
135	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	135	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
136	- L0A Buffer容量上限：L0A Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发异常。	136	- L0A Buffer容量上限：L0A Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发异常。
137	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。	137	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。
138		138

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0a_mx.md+1-1File content is under review, please try again later

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0b_2d_arch_3510.md+1-1

	@@ -134,7 +134,7 @@ PIPE_MTE1
134		134
135	- 本接口非AIC调用直接返回。	135	- 本接口非AIC调用直接返回。
136	- src位于L1 Buffer，dst位于L0B Buffer，二者位于不同的物理存储单元，不存在地址重叠。	136	- src位于L1 Buffer，dst位于L0B Buffer，二者位于不同的物理存储单元，不存在地址重叠。
137	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	137	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
138	- L0B Buffer容量上限：L0B Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发写溢出异常。	138	- L0B Buffer容量上限：L0B Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发写溢出异常。
139	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。	139	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。
140		140

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0b_3d_arch_3510.md+1-1

	@@ -126,7 +126,7 @@ PIPE_MTE1
126	- 本接口非AIC调用直接返回。	126	- 本接口非AIC调用直接返回。
127	- dst起始地址需512字节对齐（L0B Buffer对齐要求），否则会导致搬运异常。	127	- dst起始地址需512字节对齐（L0B Buffer对齐要求），否则会导致搬运异常。
128	- src起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	128	- src起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
129	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	129	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
130	- L0B Buffer容量上限：L0B Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发异常。	130	- L0B Buffer容量上限：L0B Buffer总容量64KB，dst偏移量与搬运大小之和不可越界，否则触发异常。
131	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。	131	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与源矩阵占用大小之和不可越界，否则触发异常。
132		132

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0b_mx.md+1-1

	@@ -86,7 +86,7 @@ PIPE_MTE1
86		86
87	- 本接口非AIC调用直接返回。	87	- 本接口非AIC调用直接返回。
88	- `src`位于L1 Buffer，dst位于L0B_MX Buffer，二者位于不同的物理存储单元，不存在地址重叠。	88	- `src`位于L1 Buffer，dst位于L0B_MX Buffer，二者位于不同的物理存储单元，不存在地址重叠。
89	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	89	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
90	- L0B_MX Buffer容量上限：L0B_MX Buffer总容量4KB，`dst`偏移量与搬运大小之和不可越界，否则触发地址溢出异常。	90	- L0B_MX Buffer容量上限：L0B_MX Buffer总容量4KB，`dst`偏移量与搬运大小之和不可越界，否则触发地址溢出异常。
91	- L1 Buffer容量上限：L1 Buffer总容量512KB，`src`偏移量与源矩阵占用大小之和不可越界，否则触发地址溢出异常。	91	- L1 Buffer容量上限：L1 Buffer总容量512KB，`src`偏移量与源矩阵占用大小之和不可越界，否则触发地址溢出异常。
92		92

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0b_trans_arch_3510.md+1-1

	@@ -154,7 +154,7 @@ PIPE_MTE1
154	- 本接口非AIC调用直接返回。	154	- 本接口非AIC调用直接返回。
155	- src起始地址需32字节对齐（L1 Buffer对齐要求），否则触发地址对齐异常。	155	- src起始地址需32字节对齐（L1 Buffer对齐要求），否则触发地址对齐异常。
156	- dst起始地址需512字节对齐（L0B Buffer对齐要求），否则触发地址对齐异常。	156	- dst起始地址需512字节对齐（L0B Buffer对齐要求），否则触发地址对齐异常。
157	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	157	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
158	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移与搬运大小之和不可越界，否则触发地址溢出异常。	158	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移与搬运大小之和不可越界，否则触发地址溢出异常。
159	- L0B Buffer容量上限：L0B Buffer总容量64KB，dst偏移与搬运大小之和不可越界，否则触发地址溢出异常。	159	- L0B Buffer容量上限：L0B Buffer总容量64KB，dst偏移与搬运大小之和不可越界，否则触发地址溢出异常。
160	- `repeat`为0时表示不执行搬运，本接口被视为NOP（空操作）。	160	- `repeat`为0时表示不执行搬运，本接口被视为NOP（空操作）。

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load/asc_copy_l12l0c.md+1-1

	@@ -77,7 +77,7 @@ PIPE_MTE1
77	## 约束说明	77	## 约束说明
78		78
79	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	79	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
80	-- 如果需要执行多条asc_copy_l12l0c指令，且asc_copy_l12l0c指令的目的地址存在重叠，需要插入同步指令，保证多个asc_copy_l12l0c指令的串行化，防止出现异常数据。	80	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
81		81
82	## 调用示例	82	## 调用示例
83		83

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_store/asc_copy_l0c2gm_arch_3510.md+1-1

	@@ -160,7 +160,7 @@ PIPE_FIX
160	- 本接口非AIC调用直接返回。	160	- 本接口非AIC调用直接返回。
161	- `src`起始地址需64字节对齐，否则会导致搬运异常。	161	- `src`起始地址需64字节对齐，否则会导致搬运异常。
162	- `dst`起始地址需1字节对齐，否则会导致搬运异常。	162	- `dst`起始地址需1字节对齐，否则会导致搬运异常。
163	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令，详见[核内同步](../../sync/system_sync_overview.md)。	163	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
164	- L0C Buffer容量上限：L0C Buffer总容量256KB，src偏移量与搬运大小之和不可超过L0C Buffer容量，否则触发源地址越界异常。	164	- L0C Buffer容量上限：L0C Buffer总容量256KB，src偏移量与搬运大小之和不可超过L0C Buffer容量，否则触发源地址越界异常。
165	- Nz矩阵以16×16个元素为一个基本分形。边界分形中超出`n_size`或`m_size`指定范围的数据不属于有效搬运结果。	165	- Nz矩阵以16×16个元素为一个基本分形。边界分形中超出`n_size`或`m_size`指定范围的数据不属于有效搬运结果。
166		166

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_store/asc_copy_l0c2l1_arch_2201.md+1-1

	@@ -90,7 +90,7 @@ PIPE_FIX
90		90
91	- 量化参数不能为inf/nan和非规格化数。	91	- 量化参数不能为inf/nan和非规格化数。
92	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	92	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
93	-- 如果需要执行多条asc_copy_l0c2l1指令，且asc_copy_l0c2l1指令的目的地址存在重叠，需要插入同步指令，保证多个asc_copy_l0c2l1指令的串行化，防止出现异常数据。	93	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
94	- 量化与激活模式中使用的量化系数不可为INF/NAN和非规格化数，否则会导致量化激活结果错误。	94	- 量化与激活模式中使用的量化系数不可为INF/NAN和非规格化数，否则会导致量化激活结果错误。
95		95
96	## 调用示例	96	## 调用示例

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_store/asc_copy_l0c2l1_arch_3510.md+1-1

	@@ -158,7 +158,7 @@ PIPE_FIX
158	- 本接口仅在AIC上生效，非AIC调用直接返回。	158	- 本接口仅在AIC上生效，非AIC调用直接返回。
159	- `dst`起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	159	- `dst`起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
160	- `src`起始地址需64字节对齐（L0C Buffer对齐要求，`src`为`int32_t`或`float`时），否则会导致搬运异常。	160	- `src`起始地址需64字节对齐（L0C Buffer对齐要求，`src`为`int32_t`或`float`时），否则会导致搬运异常。
161	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	161	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
162	- L0C Buffer容量上限：L0C Buffer总容量256KB，`src`偏移加搬运大小不可越界，否则触发异常。	162	- L0C Buffer容量上限：L0C Buffer总容量256KB，`src`偏移加搬运大小不可越界，否则触发异常。
163	- L1 Buffer容量上限：L1 Buffer总容量512KB，`dst`偏移加搬运大小不可越界，否则触发异常。	163	- L1 Buffer容量上限：L1 Buffer总容量512KB，`dst`偏移加搬运大小不可越界，否则触发异常。
164		164

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_store/asc_copy_l0c2ub.md+1-1

	@@ -164,7 +164,7 @@ PIPE_FIX
164	- 本接口非AIC调用直接返回。	164	- 本接口非AIC调用直接返回。
165	- `dst`起始地址需要按照32字节对齐，否则触发异常。	165	- `dst`起始地址需要按照32字节对齐，否则触发异常。
166	- `src`起始地址需要按照64字节对齐，否则触发异常。	166	- `src`起始地址需要按照64字节对齐，否则触发异常。
167	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	167	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
168	- UB容量上限：UB总容量为256KB，默认预留6KB SIMD VF栈与2KB Ascend C预留空间后可用248KB；SIMD与SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少。dst偏移量与搬运大小之和不可超过实际可用容量，否则触发写溢出异常。	168	- UB容量上限：UB总容量为256KB，默认预留6KB SIMD VF栈与2KB Ascend C预留空间后可用248KB；SIMD与SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少。dst偏移量与搬运大小之和不可超过实际可用容量，否则触发写溢出异常。
169	- L0C Buffer容量上限：L0C Buffer总容量256KB，src偏移量与搬运大小之和不可超过L0C Buffer容量，否则触发读溢出异常。	169	- L0C Buffer容量上限：L0C Buffer总容量256KB，src偏移量与搬运大小之和不可超过L0C Buffer容量，否则触发读溢出异常。
170	- Nz矩阵以16×16个元素为一个基本分形。边界分形中超出`n_size`或`m_size`指定范围的数据不属于有效搬运结果。	170	- Nz矩阵以16×16个元素为一个基本分形。边界分形中超出`n_size`或`m_size`指定范围的数据不属于有效搬运结果。

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_store/asc_copy_l12gm.md+1-1

	@@ -62,7 +62,7 @@ PIPE_MTE3
62	## 约束说明	62	## 约束说明
63		63
64	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	64	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
65	-- 如果需要执行多条asc_copy_l12gm指令，且asc_copy_l12gm指令的目的地址存在重叠，需要插入同步指令，保证多个asc_copy_l12gm指令的串行化，防止出现异常数据。	65	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
66		66
67	## 调用示例	67	## 调用示例
68		68

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_store/asc_copy_l12ub.md+1-1

	@@ -65,7 +65,7 @@ PIPE_MTE1
65	- 本接口非AIC调用直接返回。	65	- 本接口非AIC调用直接返回。
66	- dst_addr起始地址需32字节对齐（UB对齐要求），否则会导致搬运异常。	66	- dst_addr起始地址需32字节对齐（UB对齐要求），否则会导致搬运异常。
67	- src_addr起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。	67	- src_addr起始地址需32字节对齐（L1 Buffer对齐要求），否则会导致搬运异常。
68	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	68	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
69	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与搬运大小之和不可超过L1 Buffer容量，否则触发读溢出异常。	69	- L1 Buffer容量上限：L1 Buffer总容量512KB，src偏移量与搬运大小之和不可超过L1 Buffer容量，否则触发读溢出异常。
70	- UB容量上限：UB总容量为256KB，默认预留6KB SIMD VF栈与2KB Ascend C预留空间后可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少。dst偏移量与搬运大小之和不可超过实际可用容量，否则触发写溢出异常。	70	- UB容量上限：UB总容量为256KB，默认预留6KB SIMD VF栈与2KB Ascend C预留空间后可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少。dst偏移量与搬运大小之和不可超过实际可用容量，否则触发写溢出异常。
71	- `burst_count`、`burst_len`、`src_gap`、`dst_gap`取值需满足参数说明中取值范围，不满足导致搬运结果不符合预期。	71	- `burst_count`、`burst_len`、`src_gap`、`dst_gap`取值需满足参数说明中取值范围，不满足导致搬运结果不符合预期。

Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_load_aux_config/asc_fill_l1.md+1-1

	@@ -94,7 +94,7 @@ PIPE_MTE2
94	- 本接口非AIC调用直接返回。	94	- 本接口非AIC调用直接返回。
95	- `dst`起始地址需要按照32字节对齐（L1 Buffer对齐要求），否则触发地址对齐异常。	95	- `dst`起始地址需要按照32字节对齐（L1 Buffer对齐要求），否则触发地址对齐异常。
96	- L1 Buffer总容量为512KB。`dst`偏移量与实际占用空间之和不可超过L1 Buffer容量，否则触发地址溢出异常。	96	- L1 Buffer总容量为512KB。`dst`偏移量与实际占用空间之和不可超过L1 Buffer容量，否则触发地址溢出异常。
97	-- 如果本指令与其他指令存在目的地址重叠，需要插入同步指令（[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)），保证多个指令串行化，防止出现异常数据。	97	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
98	- `value`的数据类型需与`dst`的数据类型满足函数原型章节中的组合关系。两者数据类型不相同时，需要根据`dst`的数据类型转换`value`的值，具体请参考调用示例。	98	- `value`的数据类型需与`dst`的数据类型满足函数原型章节中的组合关系。两者数据类型不相同时，需要根据`dst`的数据类型转换`value`的值，具体请参考调用示例。
99		99
100	## 调用示例	100	## 调用示例

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_load.md+1-1

	@@ -82,7 +82,7 @@ __simd_callee__ inline void asc_load(vector_float& dst,
82	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	82	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
83	- src起始地址必须按dtype对齐，且必须在UB地址空间内且不越界，否则会报错。	83	- src起始地址必须按dtype对齐，且必须在UB地址空间内且不越界，否则会报错。
84	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。src起始地址加VL字节不可超过实际可用容量，否则会报错。	84	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。src起始地址加VL字节不可超过实际可用容量，否则会报错。
85	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	85	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
86		86
87	## 调用示例	87	## 调用示例
88		88

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign.md+1-1

	@@ -209,7 +209,7 @@ __simd_callee__ inline void asc_loadalign(vector_float& dst,
209	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	209	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
210	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。非连续对齐搬入模式中，当某个`DataBlock`的`mask`包含的元素全部无效时，该`DataBlock`即使越界也不会报错。	210	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。非连续对齐搬入模式中，当某个`DataBlock`的`mask`包含的元素全部无效时，该`DataBlock`即使越界也不会报错。
211	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	211	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
212	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	212	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
213		213
214	### 地址寄存器偏移搬入模式	214	### 地址寄存器偏移搬入模式
215		215

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_datablock.md+1-1

	@@ -149,7 +149,7 @@ __simd_callee__ inline void asc_loadalign_brc_datablock(vector_half& dst,
149	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	149	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
150	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	150	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
151	- 一次读取数据量为32字节（1个DataBlock）。UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。实际读取地址加32字节不可超过实际可用容量，否则会报错。	151	- 一次读取数据量为32字节（1个DataBlock）。UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。实际读取地址加32字节不可超过实际可用容量，否则会报错。
152	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	152	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
153		153
154	### 地址寄存器偏移搬入模式	154	### 地址寄存器偏移搬入模式
155		155

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_datablock_postupdate.md+1-1

	@@ -78,7 +78,7 @@ __simd_callee__ inline void asc_loadalign_brc_datablock_postupdate(vector_int8_t
78	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	78	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
79	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	79	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
80	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	80	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
81	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	81	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
82		82
83	## 调用示例	83	## 调用示例
84		84

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_elem.md+1-1

	@@ -153,7 +153,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem(vector_half& dst,
153	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	153	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
154	- 各功能模式下的实际读取地址必须按dtype对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	154	- 各功能模式下的实际读取地址必须按dtype对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
155	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	155	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
156	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	156	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
157		157
158	### 地址寄存器偏移搬入模式	158	### 地址寄存器偏移搬入模式
159		159

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_elem2datablock.md+1-1

	@@ -149,7 +149,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_half& dst,
149	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	149	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
150	- 各功能模式下，b16类型的实际读取地址必须按16字节对齐，b32类型的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	150	- 各功能模式下，b16类型的实际读取地址必须按16字节对齐，b32类型的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
151	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	151	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
152	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	152	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
153		153
154	### 地址寄存器偏移搬入模式	154	### 地址寄存器偏移搬入模式
155		155

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_elem2datablock_postupdate.md+1-1

	@@ -76,7 +76,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock_postupdate(vector_h
76	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	76	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
77	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须满足对应对齐要求：b16类型按16字节对齐，b32类型按32字节对齐；且实际读取范围必须在UB地址空间内且不越界，否则会报错。	77	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须满足对应对齐要求：b16类型按16字节对齐，b32类型按32字节对齐；且实际读取范围必须在UB地址空间内且不越界，否则会报错。
78	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	78	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
79	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	79	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
80		80
81	## 调用示例	81	## 调用示例
82		82

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_brc_elem_postupdate.md+1-1

	@@ -82,7 +82,7 @@ __simd_callee__ inline void asc_loadalign_brc_elem_postupdate(vector_int16_t& ds
82	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	82	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
83	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按dtype对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	83	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按dtype对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
84	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	84	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
85	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	85	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
86		86
87	## 调用示例	87	## 调用示例
88		88

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_datablock_strided.md+1-1

	@@ -79,7 +79,7 @@ __simd_callee__ inline vector_uint8_t asc_loadalign_datablock_strided(__ubuf__ u
79	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	79	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
80	- 实际读取地址必须按32字节对齐，且有效`DataBlock`的读取范围必须在UB地址空间内且不越界，否则会报错。	80	- 实际读取地址必须按32字节对齐，且有效`DataBlock`的读取范围必须在UB地址空间内且不越界，否则会报错。
81	- 当一个`DataBlock`中的元素全部被`mask`设置为无效时，该`DataBlock`即使越界也不会报错。	81	- 当一个`DataBlock`中的元素全部被`mask`设置为无效时，该`DataBlock`即使越界也不会报错。
82	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	82	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
83		83
84	## 调用示例	84	## 调用示例
85		85

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_deintlv.md+1-1

	@@ -153,7 +153,7 @@ __simd_callee__ inline void asc_loadalign_deintlv(vector_float& dst0,
153	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	153	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
154	- `dst0`与`dst1`必须为不同的矢量数据寄存器。	154	- `dst0`与`dst1`必须为不同的矢量数据寄存器。
155	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	155	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
156	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	156	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
157		157
158	### 地址寄存器偏移搬入模式	158	### 地址寄存器偏移搬入模式
159		159

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_deintlv_postupdate.md+1-1

	@@ -79,7 +79,7 @@ __simd_callee__ inline void asc_loadalign_deintlv_postupdate(vector_int16_t& dst
79	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	79	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
80	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	80	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
81	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	81	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
82	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	82	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
83		83
84	### 指令约束	84	### 指令约束
85		85

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_downsample.md+1-1

	@@ -173,7 +173,7 @@ __simd_callee__ inline void asc_loadalign_downsample(vector_half& dst,
173	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	173	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
174	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	174	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
175	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	175	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
176	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	176	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
177		177
178	### 地址寄存器偏移搬入模式	178	### 地址寄存器偏移搬入模式
179		179

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_downsample_postupdate.md+1-1

	@@ -88,7 +88,7 @@ __simd_callee__ inline void asc_loadalign_downsample_postupdate(vector_int16_t&
88	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	88	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
89	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	89	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
90	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	90	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
91	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	91	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
92		92
93	## 调用示例	93	## 调用示例
94		94

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_mask.md+1-1

	@@ -60,7 +60,7 @@ __simd_callee__ inline vector_bool asc_loadalign_mask(__ubuf__ uint32_t* src)
60		60
61	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	61	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
62	- `src`的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	62	- `src`的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
63	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	63	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
64		64
65	## 调用示例	65	## 调用示例
66		66

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_mask_downsample.md+1-1

	@@ -60,7 +60,7 @@ __simd_callee__ inline vector_bool asc_loadalign_mask_downsample(__ubuf__ uint32
60		60
61	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	61	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
62	- `src`的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	62	- `src`的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
63	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	63	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
64		64
65	## 调用示例	65	## 调用示例
66		66

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_mask_upsample.md+1-1

	@@ -60,7 +60,7 @@ __simd_callee__ inline vector_bool asc_loadalign_mask_upsample(__ubuf__ uint32_t
60		60
61	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	61	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
62	- `src`的实际读取地址必须按16字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	62	- `src`的实际读取地址必须按16字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
63	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	63	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
64		64
65	## 调用示例	65	## 调用示例
66		66

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_postupdate.md+1-1

	@@ -129,7 +129,7 @@ __simd_callee__ inline void asc_loadalign_postupdate(vector_float& dst,
129	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	129	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
130	- 各功能模式下的当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。非连续对齐搬入模式中，当某个`DataBlock`的`mask`包含的元素全部无效时，该`DataBlock`即使越界也不会报错。	130	- 各功能模式下的当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。非连续对齐搬入模式中，当某个`DataBlock`的`mask`包含的元素全部无效时，该`DataBlock`即使越界也不会报错。
131	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	131	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
132	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	132	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
133		133
134	### 非连续对齐搬入模式	134	### 非连续对齐搬入模式
135		135

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_unpack.md+1-1

	@@ -153,7 +153,7 @@ __simd_callee__ inline void asc_loadalign_unpack(vector_half& dst,
153	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	153	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
154	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	154	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
155	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	155	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
156	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	156	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
157		157
158	### 地址寄存器偏移搬入模式	158	### 地址寄存器偏移搬入模式
159		159

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_unpack4.md+1-1

	@@ -149,7 +149,7 @@ __simd_callee__ inline void asc_loadalign_unpack4(vector_int8_t& dst,
149	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	149	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
150	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	150	- 各功能模式下的实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
151	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	151	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
152	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	152	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
153		153
154	### 地址寄存器偏移搬入模式	154	### 地址寄存器偏移搬入模式
155		155

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_unpack4_postupdate.md+1-1

	@@ -76,7 +76,7 @@ __simd_callee__ inline void asc_loadalign_unpack4_postupdate(vector_int8_t& dst,
76	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	76	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
77	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	77	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
78	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	78	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
79	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	79	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
80		80
81	## 调用示例	81	## 调用示例
82		82

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_unpack_postupdate.md+1-1

	@@ -80,7 +80,7 @@ __simd_callee__ inline void asc_loadalign_unpack_postupdate(vector_int16_t& dst,
80	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	80	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
81	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	81	- 当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
82	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	82	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
83	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	83	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
84		84
85	## 调用示例	85	## 调用示例
86		86

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_upsample.md+1-1

	@@ -173,7 +173,7 @@ __simd_callee__ inline void asc_loadalign_upsample(vector_half& dst,
173	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	173	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
174	- 当`dst`为矢量数据寄存器时，各功能模式下的实际读取地址必须按32字节对齐；当`dst`为掩码寄存器时，实际读取地址必须按16字节对齐。实际读取范围必须在UB地址空间内且不越界，否则会报错。	174	- 当`dst`为矢量数据寄存器时，各功能模式下的实际读取地址必须按32字节对齐；当`dst`为掩码寄存器时，实际读取地址必须按16字节对齐。实际读取范围必须在UB地址空间内且不越界，否则会报错。
175	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	175	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
176	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	176	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
177		177
178	### 地址寄存器偏移搬入模式	178	### 地址寄存器偏移搬入模式
179		179

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_upsample_postupdate.md+1-1

	@@ -88,7 +88,7 @@ __simd_callee__ inline void asc_loadalign_upsample_postupdate(vector_half& dst,
88	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	88	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
89	- 当`dst`为矢量数据寄存器时，当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐；当`dst`为掩码寄存器时，必须按16字节对齐。实际读取范围必须在UB地址空间内且不越界，否则会报错。	89	- 当`dst`为矢量数据寄存器时，当前实际读取地址及Post Update更新后的下一次实际读取地址必须按32字节对齐；当`dst`为掩码寄存器时，必须按16字节对齐。实际读取范围必须在UB地址空间内且不越界，否则会报错。
90	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	90	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
91	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	91	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
92		92
93	## 调用示例	93	## 调用示例
94		94

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadunalign.md+1-1

	@@ -74,7 +74,7 @@ __simd_callee__ inline void asc_loadunalign(vector_int8_t& dst,
74	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	74	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
75	- `src1`起始地址必须按dtype对齐，且实际访问范围必须在UB地址空间内，否则会报错。	75	- `src1`起始地址必须按dtype对齐，且实际访问范围必须在UB地址空间内，否则会报错。
76	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	76	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
77	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	77	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
78	- 每次搬入前必须调用[asc_loadunalign_pre](asc_loadunalign_pre.md)中的源地址预处理模式进行初始化，且初始化接口和本接口必须使用相同的`src1`。	78	- 每次搬入前必须调用[asc_loadunalign_pre](asc_loadunalign_pre.md)中的源地址预处理模式进行初始化，且初始化接口和本接口必须使用相同的`src1`。
79		79
80	## 调用示例	80	## 调用示例

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadunalign_postupdate.md+1-1

	@@ -126,7 +126,7 @@ __simd_callee__ inline void asc_loadunalign_postupdate(vector_int8_t& dst,
126	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	126	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
127	- `src1`起始地址或`src1`与`offset`确定的实际访问地址必须按dtype对齐，且实际访问范围必须在UB地址空间内，否则会报错。	127	- `src1`起始地址或`src1`与`offset`确定的实际访问地址必须按dtype对齐，且实际访问范围必须在UB地址空间内，否则会报错。
128	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	128	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
129	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	129	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
130		130
131	### 连续非对齐搬入模式	131	### 连续非对齐搬入模式
132		132

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadunalign_pre.md+1-1

	@@ -115,7 +115,7 @@ __simd_callee__ inline void asc_loadunalign_pre(vector_load_unalign& dst,
115	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	115	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
116	- 实际访问地址在接口内部向低地址方向对齐到32字节边界后读取数据。传入的`src`或`src`与`offset`确定的地址必须按dtype对齐，对齐后的32字节读取范围必须在UB地址空间内，否则会报错。	116	- 实际访问地址在接口内部向低地址方向对齐到32字节边界后读取数据。传入的`src`或`src`与`offset`确定的地址必须按dtype对齐，对齐后的32字节读取范围必须在UB地址空间内，否则会报错。
117	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。	117	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈 + 2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB～128KB作Data Cache，可用容量进一步减少）。UB地址偏移后不可超过实际可用容量，否则会报错。
118	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	118	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
119		119
120	### 源地址预处理模式	120	### 源地址预处理模式
121		121

Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/reg_load_overview.md+1-1

	@@ -96,7 +96,7 @@ Reg数据搬入接口用于将Unified Buffer（UB）中的数据搬入矢量数
96	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	96	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
97	- 各功能模式下的实际读取地址必须满足[接口分类与选择](#接口分类与选择)中的地址对齐要求，且实际读取范围必须在UB地址空间内且不越界，否则会报错。	97	- 各功能模式下的实际读取地址必须满足[接口分类与选择](#接口分类与选择)中的地址对齐要求，且实际读取范围必须在UB地址空间内且不越界，否则会报错。
98	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化。UB地址偏移后不可超过实际可用容量，否则会报错。详细说明请参见[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	98	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化。UB地址偏移后不可超过实际可用容量，否则会报错。详细说明请参见[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
99	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	99	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
100		100
101	## 关键特性说明	101	## 关键特性说明
102		102

Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_mask/asc_get_mask_spr.md+36-33

	@@ -54,6 +54,7 @@ vector_bool，掩码寄存器。
54		54
55	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	55	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
56	- 本接口需在Vector Function（`__simd_vf__`标记的函数）内调用。	56	- 本接口需在Vector Function（`__simd_vf__`标记的函数）内调用。
		57	+- [asc_set_vector_mask](../../vector_compute/vector_mask_config/asc_set_vector_mask.md)用于在VF场景外设置掩码寄存器，在Reg矢量计算配置时，通过`asc_get_mask_spr_b16`或`asc_get_mask_spr_b32`配合获取掩码寄存器的值。更推荐用户使用VF场景内的接口（[asc_create_mask](./asc_create_mask.md)和[asc_update_mask.md](./asc_update_mask.md)）完成Reg矢量计算掩码配置。
57		58
58	## 调用示例	59	## 调用示例
59		60
	@@ -110,34 +111,31 @@ bool compare_range_data(const std::vector<T>& actual, const std::vector<T>& expe
110	return true;	111	return true;
111	}	112	}
112		113
113	-constexpr uint32_t BUFFER_BYTES = 768;	114	+constexpr uint16_t ONE_REPEAT_ELEMENTS = 256 / sizeof(uint16_t);
		115	+constexpr uint16_t REPEAT_TIMES = 2;
		116	+constexpr uint32_t BUFFER_ELEMENTS = ONE_REPEAT_ELEMENTS * REPEAT_TIMES;
		117	+constexpr uint32_t BUFFER_BYTES = BUFFER_ELEMENTS * sizeof(uint16_t);
114		118
115	-__simd_vf__ inline void exercise(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input)	119	+__simd_vf__ inline void get_mask_spr_vf(__ubuf__ uint16_t* output,
		120	+ uint16_t one_repeat_size, uint16_t repeat_time)
116	{	121	{
117	- uint32_t count16 = 128;
118	- vector_bool seed16 = asc_update_mask_b16(count16);
119	vector_bool mask16 = asc_get_mask_spr_b16();	122	vector_bool mask16 = asc_get_mask_spr_b16();
120	- vector_uint16_t value16;	123	+ vector_uint16_t dst_reg;
121	- asc_duplicate_scalar(value16, static_cast<uint16_t>(1), seed16);	124	+ for (uint16_t repeat_index = 0; repeat_index < repeat_time; ++repeat_index) {
122	- asc_storealign(reinterpret_cast<__ubuf__ uint16_t*>(output), value16, mask16);	125	+ asc_duplicate_scalar(dst_reg, static_cast<uint16_t>(1), mask16);
123	- uint32_t count32 = 64;	126	+ asc_storealign(output + repeat_index * one_repeat_size, dst_reg, mask16);
124	- vector_bool seed32 = asc_update_mask_b32(count32);	127	+ }
125	- vector_bool mask32 = asc_get_mask_spr_b32();
126	- vector_uint32_t value32;
127	- asc_duplicate_scalar(value32, static_cast<uint32_t>(1), seed32);
128	- asc_storealign(reinterpret_cast<__ubuf__ uint32_t*>(output + 256), value32, mask32);
129	}	128	}
130		129
131	-__global__ __vector__ void asc_get_mask_spr_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input)	130	+__global__ __vector__ void asc_get_mask_spr_kernel(__gm__ uint16_t* output, __gm__ uint16_t* input)
132	{	131	{
133	asc_init();	132	asc_init();
134	- __ubuf__ uint8_t output_local[BUFFER_BYTES];	133	+ __ubuf__ uint16_t output_local[BUFFER_ELEMENTS];
135	- __ubuf__ uint8_t input_local[BUFFER_BYTES];
136	- asc_copy_gm2ub_align(input_local, input, BUFFER_BYTES);
137	asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES);	134	asc_copy_gm2ub_align(output_local, input, BUFFER_BYTES);
138	asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);	135	asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
139	asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);	136	asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
140	- exercise(output_local, input_local);	137	+ asc_set_vector_mask(0, 0xffffffffULL);
		138	+ asc_vf_call<get_mask_spr_vf>(output_local, ONE_REPEAT_ELEMENTS, REPEAT_TIMES);
141	asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);	139	asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
142	asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);	140	asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
143	asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES);	141	asc_copy_ub2gm_align(output, output_local, BUFFER_BYTES);
	@@ -147,30 +145,35 @@ __global__ __vector__ void asc_get_mask_spr_kernel(__gm__ uint8_t* output, __gm_
147		145
148	int main()	146	int main()
149	{	147	{
150	- std::vector<uint8_t> input(BUFFER_BYTES), output(BUFFER_BYTES, 0xff), golden(BUFFER_BYTES);	148	+ std::vector<uint16_t> input(BUFFER_ELEMENTS), output(BUFFER_ELEMENTS, 0xffff), golden(BUFFER_ELEMENTS);
151	- for (uint32_t i = 0; i < BUFFER_BYTES; ++i) input[i] = static_cast<uint8_t>(i % 251);	149	+ for (uint32_t index = 0; index < BUFFER_ELEMENTS; ++index) {
152	- for (uint32_t i = 0; i < 256; i += sizeof(uint16_t)) golden[i] = 1;	150	+ input[index] = static_cast<uint16_t>(index + 2);
153	- for (uint32_t i = 256; i < 512; i += sizeof(uint32_t)) golden[i] = 1;	151	+ }
154	- std::copy(input.begin() + 512, input.end(), golden.begin() + 512);	152	+ golden = input;
		153	+ for (uint16_t repeat_index = 0; repeat_index < REPEAT_TIMES; ++repeat_index) {
		154	+ for (uint32_t index = 0; index < 32; ++index) {
		155	+ golden[repeat_index * ONE_REPEAT_ELEMENTS + index] = 1;
		156	+ }
		157	+ }
155	aclInit(nullptr);	158	aclInit(nullptr);
156	aclrtSetDevice(0);	159	aclrtSetDevice(0);
157	- uint8_t* input_device = nullptr;	160	+ uint16_t* input_device = nullptr;
158	- aclrtMalloc(reinterpret_cast<void*>(&input_device), (BUFFER_BYTES) sizeof(uint8_t),	161	+ aclrtMalloc(reinterpret_cast<void**>(&input_device), BUFFER_BYTES,
159	ACL_MEM_MALLOC_HUGE_FIRST);	162	ACL_MEM_MALLOC_HUGE_FIRST);
160	- uint8_t* output_device = nullptr;	163	+ uint16_t* output_device = nullptr;
161	- aclrtMalloc(reinterpret_cast<void*>(&output_device), (BUFFER_BYTES) sizeof(uint8_t),	164	+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BUFFER_BYTES,
162	ACL_MEM_MALLOC_HUGE_FIRST);	165	ACL_MEM_MALLOC_HUGE_FIRST);
163	- aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t),	166	+ aclrtMemcpy(input_device, BUFFER_BYTES, input.data(), BUFFER_BYTES,
164	ACL_MEMCPY_HOST_TO_DEVICE);	167	ACL_MEMCPY_HOST_TO_DEVICE);
165	asc_get_mask_spr_kernel<<<1, 0>>>(output_device, input_device);	168	asc_get_mask_spr_kernel<<<1, 0>>>(output_device, input_device);
166	aclrtSynchronizeDevice();	169	aclrtSynchronizeDevice();
167	- aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t),	170	+ aclrtMemcpy(output.data(), BUFFER_BYTES, output_device, BUFFER_BYTES,
168	ACL_MEMCPY_DEVICE_TO_HOST);	171	ACL_MEMCPY_DEVICE_TO_HOST);
169	- print_data("Input bytes", input);	172	+ print_data("Input elements", input);
170	- print_data("Output bytes", output);	173	+ print_data("Output elements", output);
171	- print_data("Golden bytes", golden);	174	+ print_data("Golden elements", golden);
172	const bool passed = compare_data(output, golden);	175	const bool passed = compare_data(output, golden);
173	- std::cout << (passed ? "[Success] asc_get_mask_spr_b16/asc_get_mask_spr_b32 passed." : "[Failed] asc_get_mask_spr_b16/asc_get_mask_spr_b32 failed.") << std::endl;	176	+ std::cout << (passed ? "[Success] asc_get_mask_spr_b16 passed." : "[Failed] asc_get_mask_spr_b16 failed.") << std::endl;
174	aclrtFree(input_device);	177	aclrtFree(input_device);
175	aclrtFree(output_device);	178	aclrtFree(output_device);
176	aclrtResetDevice(0);	179	aclrtResetDevice(0);

Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_permute_sel/asc_squeeze_and_storeunalign.md+1-1

	@@ -88,7 +88,7 @@ PIPE_V
88		88
89	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	89	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
90	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化。`dst_addr`加AR寄存器记录的累计字节数不可超过实际可用容量，否则会触发写越界异常。	90	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化。`dst_addr`加AR寄存器记录的累计字节数不可超过实际可用容量，否则会触发写越界异常。
91	-- 如果本接口与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	91	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
92		92
93	### 指令约束	93	### 指令约束
94		94

Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_permute_sel/asc_squeeze_and_storeunalign_finalize.md+1-1

	@@ -76,7 +76,7 @@ PIPE_V
76		76
77	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	77	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
78	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化。目的操作数的尾块地址范围不可超过实际可用容量，否则会触发写越界异常。	78	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化。目的操作数的尾块地址范围不可超过实际可用容量，否则会触发写越界异常。
79	-- 如果本接口与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	79	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
80		80
81	### 指令约束	81	### 指令约束
82		82

Mdocs/zh/api/SIMD-API/c_api/reg_compute/scatter/asc_scatter.md+1-1

	@@ -98,7 +98,7 @@ __simd_callee__ inline void asc_scatter(__ubuf__ int8_t* dst,
98	- 源操作数和索引为矢量数据寄存器，目的操作数为UB地址。	98	- 源操作数和索引为矢量数据寄存器，目的操作数为UB地址。
99	- `dst`起始地址需32字节对齐。	99	- `dst`起始地址需32字节对齐。
100	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD与SIMT混合编程时再划分32KB～128KB作为Data Cache，可用容量进一步减少）。`index`索引后的实际写入地址不可超过实际可用容量。	100	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD与SIMT混合编程时再划分32KB～128KB作为Data Cache，可用容量进一步减少）。`index`索引后的实际写入地址不可超过实际可用容量。
101	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	101	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
102		102
103	### 分散搬出约束	103	### 分散搬出约束
104		104

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_store.md+1-1

	@@ -111,7 +111,7 @@ __simd_callee__ inline void asc_store(__ubuf__ int8_t* dst,
111	- 源操作数为矢量数据寄存器，目的操作数为UB地址，UB地址空间外的指针不可作为`dst`传入。	111	- 源操作数为矢量数据寄存器，目的操作数为UB地址，UB地址空间外的指针不可作为`dst`传入。
112	- `dst`起始地址需按照`dtype`对齐。本接口按非对齐方式将矢量数据寄存器内容写入UB，不要求对齐，不会报错。	112	- `dst`起始地址需按照`dtype`对齐。本接口按非对齐方式将矢量数据寄存器内容写入UB，不要求对齐，不会报错。
113	- UB容量上限：UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的地址不可超过实际可用容量，否则会报错。	113	- UB容量上限：UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的地址不可超过实际可用容量，否则会报错。
114	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	114	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
115		115
116	### 带搬出数量模式约束	116	### 带搬出数量模式约束
117		117

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign.md+1-1

	@@ -282,7 +282,7 @@ __simd_callee__ inline void asc_storealign(__ubuf__ uint8_t* dst,
282	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	282	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
283	- dst起始地址需32字节对齐，否则会报错。	283	- dst起始地址需32字节对齐，否则会报错。
284	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址偏移后不可超过实际可用容量，否则会报错。	284	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址偏移后不可超过实际可用容量，否则会报错。
285	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	285	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
286		286
287	### 矢量数据寄存器搬出场景	287	### 矢量数据寄存器搬出场景
288		288

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_1st.md+1-1

	@@ -151,7 +151,7 @@ __simd_callee__ inline void asc_storealign_1st(__ubuf__ int8_t* dst,
151	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	151	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
152	- 源操作数为矢量数据寄存器，目的操作数为UB地址，UB地址空间外的指针不可作为`dst`传入。	152	- 源操作数为矢量数据寄存器，目的操作数为UB地址，UB地址空间外的指针不可作为`dst`传入。
153	- UB容量上限：UB总容量256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。`dst`偏移后不可超过实际可用容量。	153	- UB容量上限：UB总容量256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。`dst`偏移后不可超过实际可用容量。
154	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	154	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
155		155
156	### 指令约束	156	### 指令约束
157		157

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_1st_postupdate.md+1-1

	@@ -81,7 +81,7 @@ __simd_callee__ inline void asc_storealign_1st_postupdate(__ubuf__ int8_t*& dst,
81		81
82	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	82	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
83	- UB容量上限：UB总容量为256KB，默认预留6KB SIMD VF栈与2KB Ascend C预留空间后可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少。`dst`起始地址与Post Update累计偏移量之和不可超过实际可用容量，否则会报错。	83	- UB容量上限：UB总容量为256KB，默认预留6KB SIMD VF栈与2KB Ascend C预留空间后可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少。`dst`起始地址与Post Update累计偏移量之和不可超过实际可用容量，否则会报错。
84	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	84	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
85		85
86	### 指令约束	86	### 指令约束
87		87

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_intlv.md+1-1

	@@ -155,7 +155,7 @@ __simd_callee__ inline void asc_storealign_intlv(__ubuf__ int8_t* dst,
155	- `dst`起始地址需32字节对齐，否则会报错。	155	- `dst`起始地址需32字节对齐，否则会报错。
156	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址偏移后不可超过实际可用容量，否则会报错。	156	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址偏移后不可超过实际可用容量，否则会报错。
157	- 通过`offset`参数偏移后的实际访问地址需落在UB地址范围内，且实际访问地址仍需32字节对齐，否则会报错。	157	- 通过`offset`参数偏移后的实际访问地址需落在UB地址范围内，且实际访问地址仍需32字节对齐，否则会报错。
158	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	158	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
159		159
160	## 调用示例	160	## 调用示例
161		161

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_pack.md+1-1

	@@ -255,7 +255,7 @@ __simd_callee__ inline void asc_storealign_pack(__ubuf__ uint8_t* dst,
255	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	255	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
256	- 源操作数为矢量数据寄存器或掩码寄存器，目的操作数为UB地址。UB地址空间外的指针不可作为目的地址传入。	256	- 源操作数为矢量数据寄存器或掩码寄存器，目的操作数为UB地址。UB地址空间外的指针不可作为目的地址传入。
257	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。目的操作数地址偏移后对应的UB范围不可超过实际可用容量，否则会报错。	257	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。目的操作数地址偏移后对应的UB范围不可超过实际可用容量，否则会报错。
258	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	258	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
259		259
260	### 矢量数据寄存器搬出场景	260	### 矢量数据寄存器搬出场景
261		261

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_pack_postupdate.md+1-1

	@@ -125,7 +125,7 @@ __simd_callee__ inline void asc_storealign_pack_postupdate(__ubuf__ uint8_t*& ds
125	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	125	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
126	- 源操作数为矢量数据寄存器或掩码寄存器，目的操作数为UB地址。UB地址空间外的指针不可作为目的地址传入。	126	- 源操作数为矢量数据寄存器或掩码寄存器，目的操作数为UB地址。UB地址空间外的指针不可作为目的地址传入。
127	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。目的操作数的写出范围和Post Update后的地址不可超过实际可用容量，否则会报错。	127	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。目的操作数的写出范围和Post Update后的地址不可超过实际可用容量，否则会报错。
128	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	128	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
129		129
130	### 矢量数据寄存器压缩搬出场景	130	### 矢量数据寄存器压缩搬出场景
131		131

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_pack_quarter.md+1-1

	@@ -155,7 +155,7 @@ __simd_callee__ inline void asc_storealign_pack_quarter(__ubuf__ int32_t* dst,
155	- `dst`起始地址需32字节对齐，否则会报错。	155	- `dst`起始地址需32字节对齐，否则会报错。
156	- `mask`必须为`vector_bool`类型掩码寄存器，未赋值的掩码寄存器内容不确定，会导致有效元素位置错误。	156	- `mask`必须为`vector_bool`类型掩码寄存器，未赋值的掩码寄存器内容不确定，会导致有效元素位置错误。
157	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。`dst`偏移与写出区域对应的UB范围不可超过实际可用容量，否则会报错。	157	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。`dst`偏移与写出区域对应的UB范围不可超过实际可用容量，否则会报错。
158	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	158	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
159		159
160	### int32_t偏移压缩搬出场景	160	### int32_t偏移压缩搬出场景
161		161

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_pack_quarter_postupdate.md+1-1

	@@ -79,7 +79,7 @@ __simd_callee__ inline void asc_storealign_pack_quarter_postupdate(__ubuf__ int3
79	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	79	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
80	- `dst`起始地址需32字节对齐。	80	- `dst`起始地址需32字节对齐。
81	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。目的操作数的写出范围和Post Update后的地址不可超过实际可用容量，否则会报错。	81	- UB总容量为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈和2KB Ascend C预留空间，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作为Data Cache，可用容量进一步减少）。目的操作数的写出范围和Post Update后的地址不可超过实际可用容量，否则会报错。
82	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	82	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
83		83
84	### Quarter Pack压缩与Post Update约束	84	### Quarter Pack压缩与Post Update约束
85		85

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_postupdate.md+1-1

	@@ -160,7 +160,7 @@ __simd_callee__ inline void asc_storealign_postupdate(__ubuf__ uint8_t*& dst,
160	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	160	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
161	- dst起始地址需32字节对齐，否则会报错。	161	- dst起始地址需32字节对齐，否则会报错。
162	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数的实际访问地址和Post Update后的地址不可超过实际可用容量，否则会报错。	162	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数的实际访问地址和Post Update后的地址不可超过实际可用容量，否则会报错。
163	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	163	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
164		164
165	### 矢量数据寄存器搬出场景	165	### 矢量数据寄存器搬出场景
166		166

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_squeeze_status.md+1-1

	@@ -57,7 +57,7 @@ __simd_callee__ inline void asc_storealign_squeeze_status(__ubuf__ uint32_t* dst
57	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	57	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
58	- 实际目的地址需4字节对齐，且不得超出实际可用UB范围；无`offset`参数时实际目的地址为`dst`，带`offset`参数时为`dst`加`offset`字节。	58	- 实际目的地址需4字节对齐，且不得超出实际可用UB范围；无`offset`参数时实际目的地址为`dst`，带`offset`参数时为`dst`加`offset`字节。
59	- 调用本接口前，需使用[asc_squeeze_and_storeunalign](../reg_permute_sel/asc_squeeze_and_storeunalign.md)完成squeeze操作；写入的值是其当前累计的有效数据字节数。	59	- 调用本接口前，需使用[asc_squeeze_and_storeunalign](../reg_permute_sel/asc_squeeze_and_storeunalign.md)完成squeeze操作；写入的值是其当前累计的有效数据字节数。
60	-- 如果本接口与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	60	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
61		61
62	## 调用示例	62	## 调用示例
63		63

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storealign_squeeze_status_postupdate.md+1-1

	@@ -55,7 +55,7 @@ __simd_callee__ inline void asc_storealign_squeeze_status_postupdate(__ubuf__ ui
55	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	55	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
56	- 初始目的地址及每次Post Update后的目的地址均需4字节对齐，且不得超出实际可用UB范围。	56	- 初始目的地址及每次Post Update后的目的地址均需4字节对齐，且不得超出实际可用UB范围。
57	- 调用本接口前，需使用[asc_squeeze_and_storeunalign](../reg_permute_sel/asc_squeeze_and_storeunalign.md)完成squeeze操作；写入的值是其当前累计的有效数据字节数。	57	- 调用本接口前，需使用[asc_squeeze_and_storeunalign](../reg_permute_sel/asc_squeeze_and_storeunalign.md)完成squeeze操作；写入的值是其当前累计的有效数据字节数。
58	-- 如果本接口与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化。	58	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
59		59
60	## 调用示例	60	## 调用示例
61		61

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storeunalign.md+1-1

	@@ -86,7 +86,7 @@ __simd_callee__ inline void asc_storeunalign(__ubuf__ uint8_t* dst,
86		86
87	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	87	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
88	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址偏移后不可超过实际可用容量，否则会报错。	88	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址偏移后不可超过实际可用容量，否则会报错。
89	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	89	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
90		90
91	### 指令约束	91	### 指令约束
92		92

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storeunalign_post.md+1-1

	@@ -115,7 +115,7 @@ __simd_callee__ inline void asc_storeunalign_post(__ubuf__ uint8_t* dst,
115	- 需要保证目的操作数的地址加上`offset`对应的偏移地址，访问范围须位于实际可用UB范围内。	115	- 需要保证目的操作数的地址加上`offset`对应的偏移地址，访问范围须位于实际可用UB范围内。
116	- 该接口中的目的地址不需要32B对齐，但数据类型为`dtype`的`dst`需要`sizeof(dtype)`字节对齐。	116	- 该接口中的目的地址不需要32B对齐，但数据类型为`dtype`的`dst`需要`sizeof(dtype)`字节对齐。
117	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址不可超过实际可用容量，否则会报错。	117	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址不可超过实际可用容量，否则会报错。
118	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	118	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
119		119
120	### 立即数偏移搬出模式	120	### 立即数偏移搬出模式
121		121

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storeunalign_post_postupdate.md+1-1

	@@ -79,7 +79,7 @@ __simd_callee__ inline void asc_storeunalign_post_postupdate(__ubuf__ uint8_t*&
79	- 调用本接口前，须调用[asc_storeunalign_postupdate](asc_storeunalign_postupdate.md)的立即数偏移搬出模式，并复用同一个`unalign_reg`。	79	- 调用本接口前，须调用[asc_storeunalign_postupdate](asc_storeunalign_postupdate.md)的立即数偏移搬出模式，并复用同一个`unalign_reg`。
80	- `dst`无需32字节对齐，但须按`sizeof(dtype)`字节对齐。	80	- `dst`无需32字节对齐，但须按`sizeof(dtype)`字节对齐。
81	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址更新后不可超过实际可用容量，否则会报错。	81	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址更新后不可超过实际可用容量，否则会报错。
82	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	82	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
83		83
84	### 指令约束	84	### 指令约束
85		85

Mdocs/zh/api/SIMD-API/c_api/reg_compute/store/asc_storeunalign_postupdate.md+1-1

	@@ -165,7 +165,7 @@ __simd_callee__ inline void asc_storeunalign_postupdate(__ubuf__ uint16_t*& dst,
165		165
166	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。	166	- Reg矢量计算C API通用约束请参见[通用约束](../overview.md#通用约束)。
167	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址更新后不可超过实际可用容量，否则会报错。	167	- UB容量上限为256KB，用户可用容量随编译选项与编程场景变化（默认预留6KB SIMD VF栈+2KB Ascend C预留，可用248KB；SIMD+SIMT混编时再划分32KB~128KB作Data Cache，可用容量进一步减少）。目的操作数地址更新后不可超过实际可用容量，否则会报错。
168	-- 如果本指令与其他指令存在UB地址重叠，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。	168	+- 如果本指令与其他指令存在数据依赖，需要插入同步指令[asc_mem_bar](../reg_sync/asc_mem_bar.md)，保证多个指令串行化，防止出现异常数据。
169		169
170	### 指令约束	170	### 指令约束
171		171

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_gm2ub/asc_copy_gm2ub_arch_2201.md+1-1

	@@ -69,7 +69,7 @@ PIPE_MTE2
69	## 约束说明	69	## 约束说明
70		70
71	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	71	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
72	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	72	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
73	- 当采用前n个数据搬运接口时，搬运数据大小要求32字节对齐。	73	- 当采用前n个数据搬运接口时，搬运数据大小要求32字节对齐。
74		74
75	## 调用示例	75	## 调用示例

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_gm2ub/asc_copy_gm2ub_arch_3510.md+1-1

	@@ -104,7 +104,7 @@ PIPE_MTE2
104		104
105	- 本接口在非AIV上调用直接返回。	105	- 本接口在非AIV上调用直接返回。
106	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	106	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
107	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	107	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
108		108
109	### 连续数据搬运约束	109	### 连续数据搬运约束
110		110

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_gm2ub_align/asc_copy_gm2ub_align_arch_2201.md+2-2

	@@ -28,7 +28,7 @@
28		28
29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。	29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。
30		30
31	-提供数据非对齐搬运的功能，将数据从Global Memory (GM)搬运到Unified Buffer（UB），并支持8位/16位/32位数据类型搬运。	31	+将数据从Global Memory（GM）搬运到Unified Buffer（UB），支持搬运长度不是32字节整数倍的数据，并在搬运过程中进行32字节对齐补齐。该接口支持8位/16位/32位数据类型搬运，搬运数据长度仍需满足对应数据类型的字节对齐要求；源、目的起始地址的对齐要求请参见[参数说明](#参数说明)。
32		32
33	本接口支持以下两种数据搬运方式：	33	本接口支持以下两种数据搬运方式：
34		34
	@@ -107,7 +107,7 @@ PIPE_MTE2
107	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	107	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
108	- 当`burst_count`、`burst_len`中任意一个值为0时，该接口被视为NOP（空操作）。	108	- 当`burst_count`、`burst_len`中任意一个值为0时，该接口被视为NOP（空操作）。
109	- 当`size`值为0时，该接口被视为NOP（空操作）。	109	- 当`size`值为0时，该接口被视为NOP（空操作）。
110	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	110	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
111	- `left_padding_num`、`right_padding_num`对应的填充数据大小均不能超过32字节。	111	- `left_padding_num`、`right_padding_num`对应的填充数据大小均不能超过32字节。
112		112
113	## 调用示例	113	## 调用示例

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_gm2ub_align/asc_copy_gm2ub_align_arch_3510.md+2-2

	@@ -28,7 +28,7 @@
28		28
29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。	29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。
30		30
31	-提供数据非对齐搬运的功能，将数据从Global Memory (GM)搬运到Unified Buffer（UB），并支持8位/16位/32位数据类型搬运。	31	+将数据从Global Memory（GM）搬运到Unified Buffer（UB），支持搬运长度不是32字节整数倍的数据，并在搬运过程中进行32字节对齐补齐。该接口支持8位/16位/32位数据类型搬运，搬运数据长度仍需满足对应数据类型的字节对齐要求；源、目的起始地址的对齐要求请参见[参数说明](#参数说明)。
32		32
33	本接口支持以下两种数据搬运方式：	33	本接口支持以下两种数据搬运方式：
34		34
	@@ -171,7 +171,7 @@ PIPE_MTE2
171		171
172	- 本接口在非AIV上调用直接返回。	172	- 本接口在非AIV上调用直接返回。
173	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	173	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
174	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	174	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
175		175
176	### 连续数据搬运约束	176	### 连续数据搬运约束
177		177

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_ub2gm/asc_copy_ub2gm_arch_2201.md+1-1

	@@ -69,7 +69,7 @@ PIPE_MTE3
69	## 约束说明	69	## 约束说明
70		70
71	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	71	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
72	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	72	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
73	- 当采用前n个数据搬运接口时，搬运数据大小要求32字节对齐。	73	- 当采用前n个数据搬运接口时，搬运数据大小要求32字节对齐。
74		74
75	## 调用示例	75	## 调用示例

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_ub2gm/asc_copy_ub2gm_arch_3510.md+1-1

	@@ -102,7 +102,7 @@ PIPE_MTE3
102		102
103	- 本接口在非AIV上调用直接返回。	103	- 本接口在非AIV上调用直接返回。
104	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	104	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
105	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	105	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
106		106
107	### 连续数据搬运约束	107	### 连续数据搬运约束
108		108

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_ub2gm_align/asc_copy_ub2gm_align_arch_2201.md+2-2

	@@ -28,7 +28,7 @@
28		28
29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。	29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。
30		30
31	-将数据从Unified Buffer（UB）搬运到Global Memory (GM)，支持8位/16位/32位分块拷贝操作。	31	+将数据从Unified Buffer（UB）搬运到Global Memory（GM），支持搬运长度不是32字节整数倍的数据。搬运过程中按32字节的整数倍进行补齐处理，写入GM时丢弃补齐数据，仅写入指定长度的有效数据。该接口支持8位/16位/32位数据类型搬运，搬运数据长度仍需满足对应数据类型的字节对齐要求；源、目的起始地址的对齐要求请参见[参数说明](#参数说明)。
32		32
33	本接口支持以下两种数据搬运方式，均不支持设置填充值：	33	本接口支持以下两种数据搬运方式，均不支持设置填充值：
34		34
	@@ -105,7 +105,7 @@ PIPE_MTE3
105	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	105	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
106	- 当`burst_count`、`burst_len`中任意一个值为0时，该接口被视为NOP（空操作）。	106	- 当`burst_count`、`burst_len`中任意一个值为0时，该接口被视为NOP（空操作）。
107	- 当`size`值为0时，该接口被视为NOP（空操作）。	107	- 当`size`值为0时，该接口被视为NOP（空操作）。
108	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	108	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
109		109
110	## 调用示例	110	## 调用示例
111		111

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_ub2gm_align/asc_copy_ub2gm_align_arch_3510.md+2-2

	@@ -28,7 +28,7 @@
28		28
29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。	29	头文件路径为：`"c_api/vector_datamove/vector_datamove.h"`。
30		30
31	-提供数据非对齐搬运的功能，将数据从Unified Buffer（UB）搬运到Global Memory (GM)，并支持8位/16位/32位数据类型搬运。	31	+将数据从Unified Buffer（UB）搬运到Global Memory（GM），支持搬运长度不是32字节整数倍的数据。搬运过程中按32字节的整数倍进行补齐处理，写入GM时丢弃补齐数据，仅写入指定长度的有效数据。该接口支持8位/16位/32位数据类型搬运，搬运数据长度仍需满足对应数据类型的字节对齐要求；源、目的起始地址的对齐要求请参见[参数说明](#参数说明)。
32		32
33	本接口支持以下两种数据搬运方式，均不支持设置填充值：	33	本接口支持以下两种数据搬运方式，均不支持设置填充值：
34		34
	@@ -153,7 +153,7 @@ PIPE_MTE3
153		153
154	- 本接口在非AIV上调用直接返回。	154	- 本接口在非AIV上调用直接返回。
155	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。	155	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../../general_description_and_constraints.md#存储单元说明)。
156	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	156	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
157		157
158	### 连续数据搬运约束	158	### 连续数据搬运约束
159		159

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_ub2l1.md+1-1

	@@ -108,7 +108,7 @@ PIPE_MTE3
108		108
109	- 本接口在非AIV上调用直接返回。	109	- 本接口在非AIV上调用直接返回。
110	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../general_description_and_constraints.md#存储单元说明)。	110	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../general_description_and_constraints.md#存储单元说明)。
111	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	111	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
112		112
113	### 连续搬运模式约束	113	### 连续搬运模式约束
114		114

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_copy_ub2ub.md+1-1

	@@ -100,7 +100,7 @@ PIPE_V
100	- Ascend 950PR&950DT系列产品的`PIPE_V`内部同步由硬件自动保证，多条本接口指令之间无需额外插入同步。	100	- Ascend 950PR&950DT系列产品的`PIPE_V`内部同步由硬件自动保证，多条本接口指令之间无需额外插入同步。
101	<!-- end id9 -->	101	<!-- end id9 -->
102	<!-- npu="A3,910b" id10 -->	102	<!-- npu="A3,910b" id10 -->
103	-- Atlas A2系列产品和Atlas A3系列产品执行多条本接口指令且目的地址存在重叠时，使用`asc_sync_pipe(PIPE_V)`保证多个指令串行化。	103	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
104	<!-- end id10 -->	104	<!-- end id10 -->
105		105
106	### 连续搬运模式约束	106	### 连续搬运模式约束

Mdocs/zh/api/SIMD-API/c_api/vector_datamove/asc_ndim_copy_gm2ub.md+1-1

	@@ -137,7 +137,7 @@ PIPE_MTE2
137		137
138	- 本接口在非AIV上调用直接返回。	138	- 本接口在非AIV上调用直接返回。
139	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../general_description_and_constraints.md#存储单元说明)。	139	- 各存储单元的空间大小和对齐要求请参考[存储单元说明](../general_description_and_constraints.md#存储单元说明)。
140	-- 如果本指令与其他指令存在UB地址重叠，必须插入同步指令[asc_sync_notify](../sync/intra_core_sync/asc_sync_notify.md)和[asc_sync_wait](../sync/intra_core_sync/asc_sync_wait.md)，保证多个指令串行化，防止出现异常数据。	140	+- 如果本指令与其他指令存在数据依赖，需要插入[同步指令](../sync/system_sync_overview.md)，保证多个指令串行化，防止出现异常数据。
141	- 多核场景下，如果多个AI Core可能读取同一块GM地址，且该GM地址中的数据可能被其他核更新，应在调用本接口前调用[asc_ndim_copy_dci](asc_ndim_copy_dci.md)刷新NDDMA DataCache，保证多核间数据一致性。	141	- 多核场景下，如果多个AI Core可能读取同一块GM地址，且该GM地址中的数据可能被其他核更新，应在调用本接口前调用[asc_ndim_copy_dci](asc_ndim_copy_dci.md)刷新NDDMA DataCache，保证多核间数据一致性。
142	- 一条指令所能获取的所有数据的地址范围宽度不能超过40位（1TB），即：源操作数的i维大小为（loopi_lp_size + loopi_size + loopi_rp_size - 1） * loop0_src_stride，目的操作数的i维大小为：（loopi_lp_size + loopi_size + loopi_rp_size - 1） * loop0_dst_stride，所有的维度的大小加起来不超过2^40位。	142	- 一条指令所能获取的所有数据的地址范围宽度不能超过40位（1TB），即：源操作数的i维大小为（loopi_lp_size + loopi_size + loopi_rp_size - 1） * loop0_src_stride，目的操作数的i维大小为：（loopi_lp_size + loopi_size + loopi_rp_size - 1） * loop0_dst_stride，所有的维度的大小加起来不超过2^40位。
143	- 当每维的loopi_dst_stride为升序序列，则不同循环间的地址空间不能交织或重叠。	143	- 当每维的loopi_dst_stride为升序序列，则不同循环间的地址空间不能交织或重叠。