新增原子操作和同步资料-asc-devkit-AtomGit

已合并

新增原子操作和同步资料 #5049

GRJ_XIDUO创建于 8 天前从GRJ_XIDUO/asc-devkit_GRJ： 2026_08_12_sync_and_atomic 合入到 cann/asc-devkit： master

新增原子操作和同步资料 #5049

已合并

GRJ_XIDUO创建于 8 天前

从GRJ_XIDUO/asc-devkit_GRJ 2026_08_12_sync_and_atomic合入到cann/asc-devkit master

共 9 个文件变更+568-58

Adocs/zh/api/SIMD-API/c_api/figures/asc_sync_notify.png+3-0

@@ -0,0 +1,3 @@
	1	+version https://git-lfs.github.com/spec/v1
	2	+oid sha256:6a1eebfd7a31d7a3e419b0ebc5ce6993a33e38bb14d41c34a2f485eff21000d5
	3	+size 39901

Mdocs/zh/api/SIMD-API/c_api/figures/asc_sync_pipe.png+2-2

	@@ -1,3 +1,3 @@
1	version https://git-lfs.github.com/spec/v1	1	version https://git-lfs.github.com/spec/v1
2	-oid sha256:bfb179f7477f57802d141f14ac8364c3aba6d41ba85e978dc7512d09771f08c7	2	+oid sha256:d85e413787c86d14e1309fabacdc2a7afba9c453818e56140073b6ad797d61d2
3	-size 34066	3	+size 15843

Adocs/zh/api/SIMD-API/c_api/figures/software_synchronization_between_two_cores_flowchart.png+3-0

@@ -0,0 +1,3 @@
	1	+version https://git-lfs.github.com/spec/v1
	2	+oid sha256:545b03fd28b0d713b5b3ce075a2e22fe82f190b92d812476c0043649f31be174
	3	+size 55428

Mdocs/zh/api/SIMD-API/c_api/simd_atomic/asc_set_atomic_add.md+131-26

 ## 功能说明
++原子累加过程：将待搬运到GM的数据和GM上已有数据进行求和，然后将求和结果写入GM。本接口对后续目的地址为GM的数据搬运指令开启原子累加，不同产品支持的数据搬运通路请参考[约束说明](#约束说明)。
++
++接口可选择不同的函数原型来设定不同的累加数据类型。
+-特别地，针对Ascend 950PR/Ascend 950DT，不支持L1 Buffer到GM的通路。
+-<!-- end id8 -->
 ## 函数原型
++```c
 __aicore__ inline void asc_set_atomic_add_int8()
 __aicore__ inline void asc_set_atomic_add_int16()
 __aicore__ inline void asc_set_atomic_add_int32()
 ## 参数说明
 无
++
 ## 返回值说明
 无
 ## 约束说明
++- 各个产品由于硬件架构不同，支持的数据通路也不同，具体情况如下：
++    <!-- npu="950" id11 -->
++    - Ascend 950PR/Ascend 950DT，支持的数据通路为UB/L0C Buffer->GM。
++    <!-- end id11 -->
++    <!-- npu="A3" id8 -->
++    - Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持的数据通路为UB/L0C Buffer/L1 Buffer->GM。
++    <!-- end id8 -->
++    <!-- npu="910b" id9 -->
++    - Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持的数据通路为UB/L0C Buffer/L1 Buffer->GM。
++    <!-- end id9 -->
++- 本接口调用后会对后续所有目的地址为GM的搬运指令开启原子操作，可以调用[asc_set_atomic_none](./asc_set_atomic_none.md)接口关闭原子操作。
++- 该接口执行前不会自动将GM上已有数据置零。若开发者期望在原子累加前GM上的原始数据为零，则需手动清零。
++- 本接口仅对后续目的地址为GM的搬运指令（通过MTE1/MTE2/MTE3单元搬运）生效，对于标量写GM的指令（例如[asc_store_dev](../scalar_compute/asc_store_dev.md)）不生效。
++- 本接口与紧邻的后续搬运指令之间的同步由硬件保证，因此以下示例中插入的多流水同步是不必要的：
++
++    ```c
++    asc_set_atomic_add_int8();
++
++    /*
++    asc_set_atomic_add_int8与asc_copy_ub2gm之间的同步由硬件保证，因此以下同步是不必要的。
++    asc_sync_notify(PIPE_S, PIPE_MTE3, EVENT_ID0);
++    asc_sync_wait(PIPE_S, PIPE_MTE3, EVENT_ID0);
++    */
++
++    asc_copy_ub2gm(dst, src1, total_length * sizeof(int8_t));
++    // 关闭原子操作。
++    asc_set_atomic_none();
++    ```
++
++- 后续搬运指令的操作数据类型需与所选接口设置的数据类型一致。
++- 本接口不能保证后续搬运指令的执行顺序，若需保证确定性的执行顺序请参考[关键特性说明](key_features.md)。
 ## 调用示例
++将代码保存为`examples.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。
+-//total_length指参与搬运的数据总长度。dst是外部输入的int8_t类型的GM内存。
+-constexpr uint32_t total_length = 256;
+-__ubuf__ int8_t src0[total_length];
+-__ubuf__ int8_t src1[total_length];
+-asc_copy_ub2gm(dst, src0, total_length * sizeof(int8_t));
+-asc_sync_pipe(PIPE_MTE3);
+-asc_set_atomic_add_int8();
+-asc_copy_ub2gm(dst, src1, total_length * sizeof(int8_t));
+-asc_set_atomic_none();
+-```
++<!-- npu="950" id10 -->
++以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下：
++```bash
++bisheng examples.asc -o main --npu-arch=dav-3510; ./main
 ```
++<!-- end id10 -->
++
++```c
++#include <cstdint>
++#include <iostream>
++#include <vector>
++#include "c_api/asc_simd.h"
++#include "acl/acl.h"
++
++namespace {
++template <typename T>
++void PrintData(const char* label, const std::vector<T>& values)
++{
++    std::cout << label << ":";
++    const size_t count = values.size() < 8 ? values.size() : 8;
++    for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
++    if (values.size() > count) std::cout << " ...";
++    std::cout << std::endl;
++}
++
++template <typename T>
++bool CompareData(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0)
++{
++    if (actual.size() != expected.size()) return false;
++    for (size_t i = 0; i < actual.size(); ++i) {
++        if (actual[i] == expected[i]) continue;
++        const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
++        if (diff > tolerance || diff < -tolerance) return false;
++    }
++    return true;
++}
++
++template <typename T>
++bool CompareRangeData(const std::vector<T>& actual, const std::vector<T>& expected,
++    size_t begin, size_t count, double tolerance = 0.0)
++{
++    if (begin + count > actual.size() || begin + count > expected.size()) return false;
++    for (size_t i = begin; i < begin + count; ++i) {
++        if (actual[i] == expected[i]) continue;
++        const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
++        if (diff > tolerance || diff < -tolerance) return false;
++    }
++    return true;
++}
++
++
++
++constexpr uint32_t ELEMENTS = 8;
++
++__global__ __vector__ void AscAtomicAddKernel(__gm__ int64_t* output)
++{
++    asc_init();
++    __gm__ uint32_t* address = reinterpret_cast<__gm__ uint32_t*>(output);
++    asc_dcci_entire_all();
++    const uint32_t old_value = asc_atomic_add(address, 5U);
++    asc_sync_data_barrier(mem_dsb_t::DSB_ALL);
++    asc_dcci_entire_all();
++    output[1] = static_cast<int64_t>(old_value);
++    asc_sync();
++}
++} // namespace
++
++int main()
++{
++    std::vector<int64_t> input = {10, 0};
++    std::vector<int64_t> golden = {15, 10};
++    input.resize(ELEMENTS, 0);
++    golden.resize(ELEMENTS, 0);
++    std::vector<int64_t> output(ELEMENTS, -1);
++    aclInit(nullptr);
++    aclrtSetDevice(0);
++    int64_t* output_device = nullptr;
++    aclrtMalloc(reinterpret_cast<void**>(&output_device), (ELEMENTS) * sizeof(int64_t),
++        ACL_MEM_MALLOC_HUGE_FIRST);
++    aclrtMemcpy(output_device, input.size() * sizeof(int64_t), input.data(), input.size() * sizeof(int64_t),
++        ACL_MEMCPY_HOST_TO_DEVICE);
++    AscAtomicAddKernel<<<1, 0>>>(output_device);
++    aclrtSynchronizeDevice();
++    aclrtMemcpy(output.data(), output.size() * sizeof(int64_t), output_device, output.size() * sizeof(int64_t),
++        ACL_MEMCPY_DEVICE_TO_HOST);
++    PrintData("Input", input);
++    PrintData("Output", output);
++    PrintData("Golden", golden);
++    const bool passed = CompareData(output, golden);
++    std::cout << (passed ? "[Success] asc_atomic_add passed." : "[Failed] asc_atomic_add failed.") << std::endl;
++    aclrtFree(output_device);
++    aclrtResetDevice(0);
++    aclFinalize();
++    return passed ? 0 : 1;
++}
 ```

Adocs/zh/api/SIMD-API/c_api/simd_atomic/atomic_operations_overview.md+42-0

@@ -0,0 +1,42 @@
	1	+# 原子操作概述
	2	+
	3	+数据搬运随路原子操作接口用于对后续目的地址为GM的数据搬运指令开启原子操作，涉及的接口请参见[表1](#table1)。如下图1的左侧子图所示，未开启原子操作时，写入GM的数据搬运完成后，GM中原始数据将被新搬运数据完全覆盖。如图1右侧子图所示，当数据搬运随路原子操作接口被调用后，系统将为后续写入GM的数据搬运开启原子操作。此时，数据搬运完成后，GM中的最终数据由原始GM数据与新搬运数据共同决定。
	4	+
	5	+图1 数据搬运随路原子累加效果
	6	+![](../../../figures/atomic_operation_pipelined_data_movement_effect.png "数据搬运随路原子累加效果")
	7	+
	8	+表1 数据搬运随路原子操作接口<a name="table1"></a>
	9	+
	10	+\| 对应接口 \| 接口功能描述 \|
	11	+\| --- \| --- \|
	12	+\| [asc_set_atomic_add](asc_set_atomic_add.md) \| 对后续目的地址为GM的数据搬运开启原子累加。原子累加过程：将待拷贝的内容和GM已有内容进行求和，然后将求和结果写入GM。 \|
	13	+\| [asc_set_atomic_max](asc_set_atomic_max.md) \| 设置后续搬运到GM的数据是否执行原子比较：将待拷贝的内容和GM已有内容进行比较，然后将最大值写入GM。 \|
	14	+\| [asc_set_atomic_min](asc_set_atomic_min.md) \| 设置后续搬运到GM的数据是否执行原子比较：将待拷贝的内容和GM已有内容进行比较，然后将最小值写入GM。 \|
	15	+\| [asc_set_atomic_none](asc_set_atomic_none.md) \| 关闭数据搬运随路原子操作功能。 \|
	16	+\| [asc_set_store_atomic_config_v1](asc_set_store_atomic_config_v1.md) \| 设置数据搬运的原子操作配置。 \|
	17	+\| [asc_get_store_atomic_config](asc_get_store_atomic_config.md) \| 获取数据搬运的原子操作配置。 \|
	18	+
	19	+<!-- npu="950" id1 -->
	20	+针对Ascend 950PR/Ascend 950DT，新增Scalar原子操作接口，能够在指定GM地址上进行单点原子计算操作，涉及的接口请参见[表2](#table2)。对比数据搬运随路原子操作接口，Scalar原子操作接口不会影响后续向GM搬运数据的指令。
	21	+
	22	+如下图2左侧子图所示，不使用`asc_atomic_add`接口时，多个AI Core同时对同一GM地址执行累加操作会相互覆盖，操作不具备原子性，最终结果不可预期。如右侧子图所示，使用`asc_atomic_add`接口后，各AI Core的累加操作串行化执行，确保每次累加操作的原子性，最终结果符合预期。
	23	+
	24	+图2 标量原子累加效果
	25	+![](../../../figures/scalar_atomic_operation_diagram.png "标量原子累加效果")
	26	+
	27	+表2 Scalar原子操作接口<a name="table2"></a>
	28	+
	29	+\| 对应接口 \| 接口功能描述 \|
	30	+\| --- \| --- \|
	31	+\| [asc_atomic_add](../scalar_compute/asc_atomic_add.md) \| 用于在指定GM地址上进行原子加操作，将address指向的GM地址上的旧值（`old_value`）与输入标量值（value）求和，将和结果（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	32	+\| [asc_atomic_min](../scalar_compute/asc_atomic_min.md) \| 用于在指定GM地址上进行原子取最小值操作，将address指向的GM地址上的旧值（`old_value`）与输入标量值（value）做比较，将较小值（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	33	+\| [asc_atomic_max](../scalar_compute/asc_atomic_max.md) \| 用于在指定GM地址上进行原子取大操作，将address指向的GM地址上的旧值（`old_value`）与输入的标量值（value）进行比较，将较大值（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	34	+\| [asc_atomic_cas](../scalar_compute/asc_atomic_cas.md) \| 在指定GM地址上进行原子比较操作，读取address指向的GM地址上的旧值（`old_value`）与输入标量值value1进行比较：如果相等，则将输入标量值value2写入GM地址；如果不相等，则GM地址上的值保持不变。返回该地址修改前的值（`old_value`）。 \|
	35	+\| [asc_atomic_exch](../scalar_compute/asc_atomic_exch.md) \| 用于在GM内存中执行原子交换操作，读取address指向的GM地址上的旧值（`old_value`），并将输入的标量值（value）替换旧值存储回同一地址，返回该地址修改前的值（`old_value`）。 \|
	36	+\| [asc_atomic_and](../scalar_compute/asc_atomic_and.md) \| 对GM中的数据与指定数据执行原子与操作，即将val按位与到address指向的数据元素上。读取address指向的GM地址上的旧值（`old_value`），将旧值与输入标量值val进行按位与运算，将结果（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	37	+\| [asc_atomic_inc](../scalar_compute/asc_atomic_inc.md) \| 对GM中address指向的计数器执行原子递增操作，如果address上的数值大于等于指定数值val，则对address赋值为0，否则将address上数值加1，返回该地址修改前的值（`old_value`）。 \|
	38	+\| [asc_atomic_dec](../scalar_compute/asc_atomic_dec.md) \| 对GM中address指向的计数器执行原子递减操作，如果address上的数值等于0或大于指定数值val，则对address赋值为val，否则将address上数值减1，返回该地址修改前的值（`old_value`）。 \|
	39	+\| [asc_atomic_or](../scalar_compute/asc_atomic_or.md) \| 对GM中的数据与指定数据执行原子或操作，即将val按位或到address指向的数据元素上。读取address指向的GM地址上的旧值（`old_value`），将旧值与输入标量值val进行按位或运算，将结果（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	40	+\| [asc_atomic_sub](../scalar_compute/asc_atomic_sub.md) \| 对GM中的数据与指定数据执行原子减操作，即将val从address指向的数据元素上减去。读取address指向的GM地址上的旧值（`old_value`），将旧值减去输入标量值val，将结果（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	41	+\| [asc_atomic_xor](../scalar_compute/asc_atomic_xor.md) \| 对GM中的数据与指定数据执行原子异或操作，即将val按位异或到address指向的数据元素上。读取address指向的GM地址上的旧值（`old_value`），将旧值与输入标量值val进行按位异或运算，将结果（`new_value`）写回GM地址，返回该地址修改前的值（`old_value`）。 \|
	42	+<!-- end id1 -->

Adocs/zh/api/SIMD-API/c_api/simd_atomic/key_features.md+166-0

@@ -0,0 +1,166 @@
	1	+# 关键特性说明
	2	+
	3	+确定性计算是指在相同输入条件下，无论执行次数或执行环境如何变化，总能产生完全一致输出结果的计算过程。确定性计算为系统稳定性和实验可验证性提供保障。
	4	+
	5	+对下文描述中出现的接口有以下说明：
	6	+
	7	+- 对于浮点数类型的原子累加，本文以`asc_set_atomic_add_float()`为例，更多原子操作类型参考[asc_set_atomic_add](asc_set_atomic_add.md)。
	8	+- 对于原子最大和原子最小，本文以`asc_set_atomic_max_float()`和`asc_set_atomic_min_float()`为例，更多原子操作类型参考[asc_set_atomic_max](asc_set_atomic_max.md)和[asc_set_atomic_min](asc_set_atomic_min.md)。
	9	+- 对于AIV和AIC中向GM搬运数据的接口，本文分别以`asc_copy_ub2gm()`和`asc_copy_l0c2gm()`为例，更多搬运接口参考[矢量数据搬运](../vector_data_move/vector_data_move.md)和[矩阵数据搬运](../cube_data_move/cube_data_move.md)。
	10	+
	11	+## 确定性计算概述
	12	+
	13	+为引出原子操作场景下非确定性计算的问题，我们构建如下常见的确定性计算场景：首先，通过单组浮点数数据搬运完成GM的初始化；随后，启动原子累加操作；最后，经由多次数据搬运，在GM上对多组浮点数数据进行累加。具体伪代码如下：
	14	+
	15	+```text
	16	+1. 向GM搬运数据data0; // 数据搬运，覆盖GM原有随机值，期望GM数据为data0。
	17	+2. asc_set_atomic_add_float(); // 开启原子累加，后续从UB/L0C Buffer/L1 Buffer到GM的搬运均执行原子累加。
	18	+3. 向GM搬运data1; // 带随路原子操作的数据搬运，期望GM数据为data0 + data1。
	19	+4. 向GM搬运data2; // 带随路原子操作的数据搬运，期望GM数据为data0 + data1 + data2。
	20	+5. 向GM搬运data3; // 带随路原子操作的数据搬运，期望GM数据为data0 + data1 + data2 + data3。
	21	+```
	22	+
	23	+如下图1所示，开发者的预期结果：指令发射的顺序能够严格对应实际指令执行顺序，多次执行该段代码，无论执行多少次，最终GM数据均为data0 + data1 + data2 + data3，结果完全一致，实现确定性计算。
	24	+
	25	+图1 确定性计算场景，GM上数据变化过程
	26	+![](../../../figures/deterministic_computation_gm_data_change_process.png "Deterministic computation scenario, GM data change process")
	27	+
	28	+但实际情况是，若开发者不做干预，程序每次运行时这些指令的执行顺序都可能发生变化，最终导致GM数据与预期结果不一致。下面列举两种可能的指令执行顺序及其对应的执行流程。
	29	+
	30	+## 非确定性计算，结果1
	31	+
	32	+图2 非确定性计算场景1，GM上数据变化过程
	33	+![](../../../figures/non_deterministic_computation_1_gm_data_change_process.png "Non-deterministic computation scenario 1, GM data change process")
	34	+
	35	+如图2所示，该场景中指令执行流程如下：
	36	+
	37	+1. 初始状态，GM数据为：随机值；
	38	+2. 向GM搬运data0，GM数据被初始化为：data0；
	39	+3. 执行`asc_set_atomic_add_float()`，为后续搬运指令开启原子累加，GM数据为：data0；
	40	+4. 三次带随路原子操作的搬运指令乱序，实际执行顺序为“搬出data2→搬出data3→搬出data1”，最终GM上数据为：data0 + data2 + data3 + data1。
	41	+
	42	+非确定性计算的产生原因1：
	43	+
	44	+带随路原子操作的搬运指令乱序，由于浮点数加法不满足结合律，即\(a+b\)+c!=a+\(b+c\)，使得最终GM数据data0 + data2 + data3 + data1与预期的data0 + data1 + data2 + data3存在偏差。
	45	+
	46	+带随路原子操作的搬运指令乱序，导致最终结果产生偏差的前提条件有以下三条：
	47	+
	48	+- 原子操作类型为原子累加（最大值、最小值运算满足结合律）。
	49	+- 原子操作数据类型为浮点数（整数加法满足结合律）。
	50	+- 带随路原子操作的搬运指令达到3条及以上（浮点数加法满足交换律，但是不满足结合律）。
	51	+
	52	+## 非确定性计算，结果2
	53	+
	54	+图3 非确定性计算场景2，GM上数据变化过程
	55	+![](../../../figures/non_deterministic_computation_2_gm_data_change_process.png "Non-deterministic computation scenario 2, GM data change process")
	56	+
	57	+如图3所示，该场景中指令执行流程如下：
	58	+
	59	+1. 初始状态，GM数据为：随机值；
	60	+2. 执行`asc_set_atomic_add_float()`，为后续搬运指令开启原子累加，GM数据为：随机值；
	61	+3. 先后执行两次带随路原子操作的搬运指令，执行顺序为“搬出data1→搬出data2”，GM数据为：随机值 + data1 + data2；
	62	+4. 向GM搬运data0，GM上累加的结果被data0覆盖，GM数据为：data0；
	63	+5. 最后执行data3的搬运，最终GM上数据为：data0 + data3。
	64	+
	65	+非确定性计算的产生原因2：
	66	+
	67	+开启原子累加前的普通搬运指令与开启原子操作的搬运指令之间发生乱序，会导致GM上已完成原子操作的数据被data0错误覆盖，进而产生非确定性计算结果。
	68	+
	69	+此类乱序导致结果偏差无需满足任何前提条件，开发者无需再区分原子操作类型、原子操作数据类型，也无需考虑带随路原子操作的搬运指令数量是否达到3条及以上。
	70	+
	71	+## 确定性计算实现方案
	72	+
	73	+根据导致非确定性计算的两个根因，下面也从解决这两个方面描述确定性计算实现的方案。核心思想是在指令之间插入适当的同步，使每次程序运行时相关指令都按照预期确定的顺序执行，最终保证每次执行程序输出的结果都相同。具体来说包含以下两个方面：
	74	+
	75	+- 开启原子累加前的搬运指令与开启原子操作的指令之间插入同步
	76	+
	77	+ 如下伪代码所示，在指令1与2之间插入同步，能够确保开始原子操作前GM的初始值符合预期。
	78	+
	79	+- 开启原子累加操作后，多条搬运指令之间的同步
	80	+
	81	+ 指令3与4、4与5之间插入同步，能够确保浮点数累加的顺序符合预期。
	82	+
	83	+ >[!CAUTION]注意
	84	+ >开启原子操作的指令与后续搬运指令之间的同步由硬件保证。
	85	+
	86	+```text
	87	+// 整个原子累加在同一个核内执行，控制5个指令的执行顺序为“1→2→3→4→5”。
	88	+1. 向GM搬运数据data0; // 数据搬运，覆盖GM原有随机值，期望GM数据为data0。
	89	+核内同步
	90	+2. asc_set_atomic_add_float(); // 开启原子累加，后续从UB/L0C Buffer/L1 Buffer到GM的搬运均执行原子累加。
	91	+// 指令2与3间无需插入同步。
	92	+3. 向GM搬运data1; // 开启原子累加后的数据搬运，期望GM数据为data0 + data1。
	93	+核内同步
	94	+4. 向GM搬运data2; // 开启原子累加后的数据搬运，期望GM数据为data0 + data1 + data2。
	95	+核内同步
	96	+5. 向GM搬运data3; // 开启原子累加后的数据搬运，期望GM数据为data0 + data1 + data2 + data3。
	97	+```
	98	+
	99	+如下伪代码所示，当上述指令都在不同核中执行时，需要将上述的核内同步替换为核间同步。
	100	+
	101	+```text
	102	+// 整个原子累加在4个不同核中执行，控制4个核执行的顺序为“核0→核1→核2→核3”。
	103	+if (asc_get_block_idx() == 0) {
	104	+ 向GM搬运数据data0;
	105	+ 核间同步
	106	+} else if (asc_get_block_idx() == 1) {
	107	+ 核间同步
	108	+ asc_set_atomic_add_float();
	109	+ 向GM搬运data1;
	110	+ 核间同步
	111	+} else if (asc_get_block_idx() == 2) {
	112	+ 核间同步
	113	+ asc_set_atomic_add_float();
	114	+ 向GM搬运data2;
	115	+ 核间同步
	116	+} else if (asc_get_block_idx() == 3) {
	117	+ 核间同步
	118	+ asc_set_atomic_add_float();
	119	+ 向GM搬运data3;
	120	+}
	121	+```
	122	+
	123	+下面介绍如何基于硬件同步指令实现核内同步，以及如何基于软件同步方案实现核间同步。
	124	+
	125	+## 核内同步
	126	+
	127	+搬运指令和开启原子操作的指令流水类型如下表所示，当上述指令在同一个核内执行时，开发者按需插入[asc_sync_pipe](../sync/asc_sync_pipe.md)或者[asc_sync_notify](../sync/asc_sync_notify.md)与[asc_sync_wait](../sync/asc_sync_wait.md)。
	128	+
	129	+表1 原子操作确定性计算相关指令的流水类型
	130	+
	131	+\| 指令名称 \| 流水类型 \|
	132	+\| --- \| --- \|
	133	+\| asc_copy_ub2gm \| PIPE_MTE3 \|
	134	+\| asc_copy_l0c2gm \| PIPE_FIX \|
	135	+\| asc_set_atomic_add_float()/asc_set_atomic_max_float()/asc_set_atomic_min_float() \| PIPE_S \|
	136	+
	137	+## 核间同步
	138	+
	139	+由于当前未提供用于控制不同核之间执行顺序的硬件同步接口，因此确定性计算场景下的核间同步需通过软件仿真实现：通过GM中的信号量实现核间同步，先建立一对核（AIV或AIC）之间的同步，进而扩展至多个核之间的同步。
	140	+
	141	+下图展示了两个核之间如何通过GM中的信号量进行核间同步：
	142	+
	143	+- 上一个核完成数据搬运或开启原子操作后，会通过Scalar单元向核间共享的GM中的信号量写入值1，表示自己的任务已完成。上一个核中也需要插入核内同步：
	144	+ - 当上一个核内有多条搬运指令时，它们之间需要插入核内同步1。
	145	+ - Scalar单元向GM写数据之前必须保证前序所有搬运指令都已经执行完成，因此它们之间也需要插入核内同步2。
	146	+
	147	+- 当前核在执行搬运任务前，会通过Scalar单元不断读取该信号量的值。如果信号量不等于1，当前核会进入阻塞等待状态；当检测到信号量等于1时，当前核会解除阻塞，开始执行自己的数据搬运或原子操作。为确保信号量等于1之前，当前核不会执行搬运指令，需要在搬运指令之前插入核内同步3。
	148	+
	149	+图4 一对核之间软件同步方案流程图
	150	+![](../figures/software_synchronization_between_two_cores_flowchart.png "Software synchronization flowchart between two cores")
	151	+
	152	+Scalar单元访问GM上的信号量，存在两种访问方式：
	153	+
	154	+- 经过DCache访问
	155	+
	156	+ 当开发者通过`x_gm[i]`（其中`x_gm`为`__gm__`类型指针）读写GM时，需手动调用[asc_dcci](../cache_ctrl/asc_dcci.md)接口，以保证多核间的数据一致性。
	157	+
	158	+- 不通过DCache访问
	159	+
	160	+ 使用[asc_store_dev](../scalar_compute/asc_store_dev.md)和与其对应的绕过DCache读取能力直接访问GM。这种方式无需额外操作即可保证多核间数据的一致性。
	161	+
	162	+如图4所示，核间同步方案中也需要与核内同步配合使用，现将三处核内同步作用说明如下：
	163	+
	164	+- 核内同步1（可选）：核内存在多条数据搬运指令时，通过该同步保证各搬运操作严格按顺序执行。
	165	+- 核内同步2（必备）：等待前一个核全部任务执行完毕后，才允许Scalar单元向全局内存信号量写入1。
	166	+- 核内同步3（必备）：等待Scalar单元检测到信号量更新为1后，当前核再启动后续任务执行。

Mdocs/zh/api/SIMD-API/c_api/sync/asc_sync_notify.md+112-29

	@@ -26,12 +26,20 @@
26		26
27	## 功能说明	27	## 功能说明
28		28
29	-设置同步标志，通知目标流水线。	29	+如图1所示，与[asc_sync_wait](asc_sync_wait.md)配对使用，用于实现AI Core内部不同流水之间的同步控制，`asc_sync_notify`和`asc_sync_wait`各自的功能如下：
		30	+
		31	+- `asc_sync_notify`：当源流水的前序指令的所有读写操作都完成之后，当前指令开始执行，并将硬件中的对应标志位设置为1。`asc_sync_notify`只是设置硬件中的对应标志位，并不会阻塞源流水中的下一个指令。
		32	+- `asc_sync_wait`：当目的流水执行到该指令时，如果发现硬件中对应标志位为0，目的流水的后续指令将一直被阻塞；如果发现硬件中对应标志位为1，则将硬件中对应标志位设置为0，同时目的流水的后续指令开始执行。
		33	+
		34	+图1 asc_sync_notify和asc_sync_wait接口功能示意图
		35	+![](../../c_api/figures/asc_sync_notify.png)
30		36
31	## 函数原型	37	## 函数原型
32		38
33	-```cpp	39	+```c
34	-__aicore__ inline void asc_sync_notify(pipe_t pipe, pipe_t tpipe, event_t id)	40	+__aicore__ inline void asc_sync_notify(pipe_t pipe,
		41	+ pipe_t tpipe,
		42	+ event_t id)
35	```	43	```
36		44
37	## 参数说明	45	## 参数说明
	@@ -39,10 +47,25 @@ __aicore__ inline void asc_sync_notify(pipe_t pipe, pipe_t tpipe, event_t id)
39	表1 参数说明	47	表1 参数说明
40		48
41	\| 参数名 \| 输入/输出 \| 描述 \|	49	\| 参数名 \| 输入/输出 \| 描述 \|
42	-\| :--- \| :--- \| :--- \|	50	+\|---\|---\|---\|
43	-\| pipe \| 输入 \| 源流水线类型。需传入编译期常量。 \|	51	+\| pipe \| 输入 \| 源流水类型，即“等待哪条流水的前序指令完成”。取值范围为`pipe_t`枚举：`PIPE_S`、`PIPE_V`、`PIPE_M`、`PIPE_MTE1`、`PIPE_MTE2`、`PIPE_MTE3`、`PIPE_FIX`。 \|
		M M munanhw7 天前应修改为中文双引号
44	-\| tpipe \| 输入 \| 目标流水线类型。需传入编译期常量。 \|	52	+\| tpipe \| 输入 \| 目标流水类型，即“解除哪条流水的`asc_sync_wait`阻塞”。取值范围与`pipe`相同，为`pipe_t`枚举。 \|
45	-\| id \| 输入 \| 同步ID。取值范围为[0, 7]。 \|	53	+\| id \| 输入 \| 同步事件ID，每对`pipe`与`tpipe`组合各自拥有8个独立的同步事件ID。取值范围为`event_t`枚举类型。 \|
		54	+
		55	+`event_t`枚举定义如下：
		56	+
		57	+```c
		58	+typedef enum {
		59	+ EVENT_ID0 = 0,
		60	+ EVENT_ID1 = 1,
		61	+ EVENT_ID2 = 2,
		62	+ EVENT_ID3 = 3,
		63	+ EVENT_ID4 = 4,
		64	+ EVENT_ID5 = 5,
		65	+ EVENT_ID6 = 6,
		66	+ EVENT_ID7 = 7
		67	+} event_t;
		68	+```
46		69
47	## 返回值说明	70	## 返回值说明
48		71
	@@ -54,31 +77,91 @@ PIPE_S
54		77
55	## 约束说明	78	## 约束说明
56		79
57	-- asc_sync_notify和asc_sync_wait必须成对使用。	80	+- `pipe`与`tpipe`并非任意组合，两者组合的取值存在限制：针对不同产品，AIC与AIV中支持的组合不同，具体请参考[核内同步分类](intra_core_sync_overview.md#核内同步分类)中的[表2](intra_core_sync_overview.md#aic_intra_core_sync_combinations)和[表3](intra_core_sync_overview.md#aiv_intra_core_sync_combinations)。
58	-	81	+- 相同源流水、相同目标流水、相同id下，连续使用`asc_sync_notify`会引发未定义行为。
59	-- 相同源流水、相同目标流水、相同id下，连续使用asc_sync_notify会引发未定义行为。	82	+- 本接口需与`asc_sync_wait`配对使用，配对的两条调用其`pipe`、`tpipe`、`id`三个参数必须完全一致。
		83	+- `pipe`与`tpipe`均不可取`PIPE_ALL`，否则触发异常。
		84	+- 每对`pipe`与`tpipe`组合各自拥有8个独立的同步事件ID。例如`PIPE_M`与`PIPE_V`的组合和`PIPE_V`与`PIPE_MTE3`的组合可同时使用相同的`id`值而互不干扰。
		85	+- 本接口不会对 `pipe`与`tpipe`两条流水的后续指令不产生阻塞效果。
60		86
61	## 调用示例	87	## 调用示例
62		88
63	-```cpp	89	+将代码保存为`examples.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)。
64	-// 本例中total_length指参与计算的数据总长度。src0_gm，src1_gm，dst_gm是外部输入的float类型的源操作数、目的操作数，指向GM内存空间。
65	-constexpr uint32_t total_length = 128;
66	-__ubuf__ float src0[total_length];
67	-__ubuf__ float src1[total_length];
68	-__ubuf__ float dst[total_length];
69		90
70	-asc_copy_gm2ub(src0, src0_gm, total_length * sizeof(float));	91	+<!-- npu="950" id8 -->
71	-asc_copy_gm2ub(src1, src1_gm, total_length * sizeof(float));	92	+以Ascend 950PR/Ascend 950DT产品（对应NPU架构为`dav-3510`）为例，编译运行命令如下：
72		93
73	-// 同步操作：数据搬运操作（GM到UB，PIPE_MTE2流水）完成后才能启动计算操作（PIPE_V流水）。	94	+```bash
74	-asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0); // EVENT_ID0为外部传入的同步ID。	95	+bisheng examples.asc -o main --npu-arch=dav-3510; ./main
75	-asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0); // EVENT_ID0为外部传入的同步ID。	96	+```
76	-	97	+<!-- end id8 -->
77	-asc_add(dst, src1, src0, total_length);	98	+
78	-	99	+```c
79	-// 同步操作：计算操作（PIPE_V流水）完成后才能启动数据搬运操作（UB到GM，PIPE_MTE3流水）。	100	+#include <cstdint>
80	-asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0); // EVENT_ID0为外部传入的同步ID。	101	+#include <iostream>
81	-asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0); // EVENT_ID0为外部传入的同步ID。	102	+#include <vector>
82	-	103	+#include "c_api/asc_simd.h"
83	-asc_copy_ub2gm(dst_gm, dst, total_length * sizeof(float));	104	+#include "acl/acl.h"
		105	+
		106	+namespace {
		107	+
		108	+constexpr uint32_t ELEMENTS = 64;
		109	+constexpr uint32_t BYTES = ELEMENTS * sizeof(float);
		110	+
		111	+void PrintData(const char* label, const std::vector<float>& data)
		112	+{
		113	+ std::cout << label << ":";
		114	+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
		115	+ std::cout << " ..." << std::endl;
		116	+}
		117	+
		118	+__global__ __vector__ void AscSyncNotifyKernel(__gm__ float* output, __gm__ float* src0, __gm__ float* src1)
		119	+{
		120	+ asc_init();
		121	+ __ubuf__ float x[ELEMENTS], y[ELEMENTS], z[ELEMENTS];
		122	+ asc_copy_gm2ub_align(x, src0, BYTES);
		123	+ asc_copy_gm2ub_align(y, src1, BYTES);
		124	+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
		125	+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
		126	+ asc_add(z, x, y, ELEMENTS);
		127	+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
		128	+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
		129	+ asc_copy_ub2gm_align(output, z, BYTES);
		130	+ asc_sync_mte3(0);
		131	+}
		132	+
		133	+} // namespace
		134	+
		135	+int main()
		136	+{
		137	+ std::vector<float> src0(ELEMENTS), src1(ELEMENTS), output(ELEMENTS, 0.0f), golden(ELEMENTS);
		138	+ for (uint32_t i = 0; i < ELEMENTS; ++i) {
		139	+ src0[i] = static_cast<float>(i) * 0.25f;
		140	+ src1[i] = static_cast<float>(ELEMENTS - i) * 0.5f;
		141	+ golden[i] = src0[i] + src1[i];
		142	+ }
		143	+ aclInit(nullptr);
		144	+ aclrtSetDevice(0);
		145	+ float src0_device = nullptr, src1_device = nullptr, *output_device = nullptr;
		146	+ aclrtMalloc(reinterpret_cast<void**>(&src0_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
		147	+ aclrtMalloc(reinterpret_cast<void**>(&src1_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
		148	+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
		149	+ aclrtMemcpy(src0_device, BYTES, src0.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
		150	+ aclrtMemcpy(src1_device, BYTES, src1.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
		151	+ AscSyncNotifyKernel<<<1, 0>>>(output_device, src0_device, src1_device);
		152	+ aclrtSynchronizeDevice();
		153	+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
		154	+ PrintData("Input src0", src0);
		155	+ PrintData("Input src1", src1);
		156	+ PrintData("Output", output);
		157	+ PrintData("Golden", golden);
		158	+ const bool passed = output == golden;
		159	+ std::cout << (passed ? "[Success] asc_sync_notify passed." : "[Failed] asc_sync_notify failed.") << std::endl;
		160	+ aclrtFree(src0_device);
		161	+ aclrtFree(src1_device);
		162	+ aclrtFree(output_device);
		163	+ aclrtResetDevice(0);
		164	+ aclFinalize();
		165	+ return passed ? 0 : 1;
		166	+}
84	```	167	```

Adocs/zh/api/SIMD-API/c_api/sync/intra_core_sync_overview.md+108-0

@@ -0,0 +1,108 @@
	1	+# 核内同步能力概述
	2	+
	3	+## 为什么需要核内同步
	4	+
	5	+AI Core内部的执行单元（如MTE2搬运单元、Vector计算单元等）以异步并行的方式运行，在读写同一存储资源时可能存在数据依赖关系。为确保数据一致性及计算正确性，需通过同步控制协调操作时序。
	6	+
	7	+<!-- npu="950" id2 -->
	8	+
	9	+针对[NPU架构3510](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)，硬件架构图如下，高亮部分展示了并行执行的计算单元和搬运单元。
	10	+
	11	+图1 NPU架构3510架构图
	12	+![](../../../figures/ascend_950pr_950dt_architecture.png "Ascend_950PR_950DT_硬件架构图")
	13	+
	14	+<!-- end id2 -->
	15	+
	16	+<!-- npu="A3,910b" id1 -->
	17	+
	18	+针对[NPU架构2201](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)，硬件架构图如下，高亮部分展示了并行执行的计算单元和搬运单元。
	19	+
	20	+图2 NPU架构2201架构图
	21	+![](../../../figures/atlas_a2_a3_architecture.png)
	22	+
	23	+<!-- end id1 -->
	24	+
	25	+下图示例描述了一个常见的Vector计算数据流：
	26	+
	27	+1. 先通过DMA执行单元将数据从Global Memory搬入到Local Memory；
	28	+2. 进行计算；
	29	+3. 然后再通过DMA执行单元将计算结果从Local Memory搬出到Global Memory。
	30	+
	31	+图3 Vector计算数据流示意图
	32	+![](../../../figures/vector_compute_data_flow.png "Vector_计算_数据流示意图")
	33	+
	34	+四个执行单元Scalar、Vector、DMA（MTE2）、DMA（MTE3）并行执行，若访问同一片Local Memory，需要同步机制来控制它们的访问时序：保证先搬入Local Memory后再计算，计算完成后再搬出。
	35	+
	36	+图4 核内并行流水执行时序示意图
	37	+![](../../../figures/intra_core_parallel_pipeline_execution_timing.png "intra_core_parallel_pipeline_execution_timing")
	38	+
	39	+## 硬件流水类型
	40	+
	41	+AI Core内部并行的指令流水类型和解释如下所示：
	42	+
	43	+> [!NOTE]说明
	44	+>不同的硬件架构，每一种硬件流水类型包含的具体流水会有所差异，详细介绍请参考[硬件实现](../../../../guide/编程指南/高级编程/硬件实现/硬件实现.md)章节。
	45	+
	46	+表1 指令流水类型和相关说明
	47	+
	48	+\| 流水类型 \| 含义 \|
	49	+\| --- \| --- \|
	50	+\| PIPE_S \| 标量流水线，使用标量访存语句或标量计算接口访问GM、片上存储地址时为此流水 \|
	51	+\| PIPE_V \| 矢量计算流水及部分硬件架构下的L0C Buffer->UB数据搬运流水 \|
	52	+\| PIPE_M \| 矩阵计算流水 \|
	53	+\| PIPE_MTE1 \| L1 Buffer->L0A Buffer、L1 Buffer->L0B Buffer数据搬运流水 \|
	54	+\| PIPE_MTE2 \| GM->L1 Buffer、GM->UB等数据搬运流水 \|
	55	+\| PIPE_MTE3 \| UB->GM等数据搬运流水 \|
	56	+\| PIPE_FIX \| L0C Buffer->GM、L0C Buffer->L1 Buffer等数据搬运流水 \|
	57	+
	58	+## 核内同步分类
	59	+
	60	+对上述核内并行流水的同步控制分为两种：
	61	+
	62	+- 多流水同步：同一核内具有数据依赖的不同类型流水指令之间的同步。
	63	+
	64	+ - 通过[asc_sync_notify](asc_sync_notify.md)/[asc_sync_wait](asc_sync_wait.md)接口进行不同流水线间的同步控制。在`asc_sync_notify`/`asc_sync_wait`的指令中，可以指定一对指令流水（源流水与目的流水）先后执行的关系，表示两个指令流水之间完成一组“锁”机制，其作用原理为：
	65	+ - `asc_sync_notify`：当源流水的前序指令的所有读写操作都完成之后，当前指令开始执行，并将硬件中的对应标志位设置为1。
	66	+ - `asc_sync_wait`：当目的流水执行到该指令时，如果发现硬件中对应标志位为0，目的流水的后续指令将一直被阻塞；如果发现硬件中对应标志位为1，则将硬件中对应标志位设置为0，同时目的流水的后续指令开始执行。
	67	+ <!-- npu="950" id3 -->
	68	+ - Ascend 950PR/Ascend 950DT新增通过[asc_lock](asc_lock.md)/[asc_unlock](asc_unlock.md)接口进行不同流水线间的同步控制。通过Lock锁定指定流水（阻塞后续指令），再通过Unlock释放流水，来完成流水间的同步依赖。
	69	+ - `asc_lock`：根据mutex_id获取Mutex，若Mutex已被锁定，将阻塞后续指定流水指令队列，直到前序指令中对应mutex_id的Mutex被`asc_unlock`。
	70	+ - `asc_unlock`：当前流水的前置指令退出后，根据mutex_id释放对应Mutex。
	71	+ <!-- end id3 -->
	72	+ - `asc_lock`与`asc_unlock`与`asc_sync_notify`和`asc_sync_wait`相比，该组合内聚性更强、与其它流水线解耦，可简化反向同步逻辑，具体说明请参考[功能说明](asc_lock.md#功能说明)。
	73	+
	74	+- 单流水同步：同一核内具有数据依赖的相同类型流水指令之间的同步。通过[asc_sync_pipe](asc_sync_pipe.md)接口进行相同流水线间的同步控制。同一流水中虽然指令是顺序执行，但并不意味着后一条指令开始执行时前一条指令执行结束。`asc_sync_pipe`指令可以保证前序指令中所有数据读写全部完成，后序指令才开始执行。注意该接口不支持PIPE_S单流水的同步。
	75	+- 通过[asc_sync_data_barrier](asc_sync_data_barrier.md)接口阻塞后续指令的执行，直到此前已发出但尚未完成的内存访问指令全部执行完成。开发者通过`arg`参数指定屏障作用的内存范围，确保屏障前后的内存访问指令按预期顺序完成。
	76	+
	77	+<!-- npu="A3,910b" id4 -->
	78	+以[NPU架构2201](../../../../guide/编程指南/语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)为例，该硬件架构下所有合法的核内同步组合如[表2](#aic_intra_core_sync_combinations)和[表3](#aiv_intra_core_sync_combinations)所示。其中，“不涉及”表示硬件层面不存在此种同步组合，“暂无应用场景”表示存在此种同步组合，但在实际开发场景中暂不需要使用。
	79	+
	80	+<a name="aic_intra_core_sync_combinations"></a>
	81	+
	82	+表2 AIC中所有合法的核内同步组合
	83	+
	84	+\| 源流水/目的流水 \| PIPE_S \| PIPE_M \| PIPE_MTE1 \| PIPE_MTE2 \| PIPE_MTE3 \| PIPE_FIX \|
	85	+\| --- \| --- \| --- \| --- \| --- \| --- \| --- \|
	86	+\| PIPE_S \| 不涉及 \| 不涉及 \| 不涉及 \| 不涉及 \| 不涉及 \| 不涉及 \|
	87	+\| PIPE_M \| 不涉及 \| asc_sync_pipe(PIPE_M) \| asc_sync_notify(PIPE_M, PIPE_MTE1, 0)<br>asc_sync_wait(PIPE_M, PIPE_MTE1, 0) \| asc_sync_notify(PIPE_M, PIPE_MTE2, 0)<br>asc_sync_wait(PIPE_M, PIPE_MTE2, 0) \| 不涉及 \| asc_sync_notify(PIPE_M, PIPE_FIX, 0)<br>asc_sync_wait(PIPE_M, PIPE_FIX, 0) \|
	88	+\| PIPE_MTE1 \| 不涉及 \| asc_sync_notify(PIPE_MTE1, PIPE_M, 0)<br>asc_sync_wait(PIPE_MTE1, PIPE_M, 0) \| asc_sync_pipe(PIPE_MTE1) \| asc_sync_notify(PIPE_MTE1, PIPE_MTE2, 0)<br>asc_sync_wait(PIPE_MTE1, PIPE_MTE2, 0) \| asc_sync_notify(PIPE_MTE1, PIPE_MTE3, 0)<br>asc_sync_wait(PIPE_MTE1, PIPE_MTE3, 0) \| asc_sync_notify(PIPE_MTE1, PIPE_FIX, 0)<br>asc_sync_wait(PIPE_MTE1, PIPE_FIX, 0) \|
	89	+\| PIPE_MTE2 \| 不涉及 \| asc_sync_notify(PIPE_MTE2, PIPE_M, 0)<br>asc_sync_wait(PIPE_MTE2, PIPE_M, 0) \| asc_sync_notify(PIPE_MTE2, PIPE_MTE1, 0)<br>asc_sync_wait(PIPE_MTE2, PIPE_MTE1, 0) \| asc_sync_pipe(PIPE_MTE2) \| asc_sync_notify(PIPE_MTE2, PIPE_MTE3, 0)<br>asc_sync_wait(PIPE_MTE2, PIPE_MTE3, 0) \| 暂无应用场景 \|
	90	+\| PIPE_MTE3 \| 不涉及 \| 不涉及 \| asc_sync_notify(PIPE_MTE3, PIPE_MTE1, 0)<br>asc_sync_wait(PIPE_MTE3, PIPE_MTE1, 0) \| asc_sync_notify(PIPE_MTE3, PIPE_MTE2, 0)<br>asc_sync_wait(PIPE_MTE3, PIPE_MTE2, 0) \| asc_sync_pipe(PIPE_MTE3) \| 暂无应用场景 \|
	91	+\| PIPE_FIX \| 不涉及 \| asc_sync_notify(PIPE_FIX, PIPE_M, 0)<br>asc_sync_wait(PIPE_FIX, PIPE_M, 0) \| asc_sync_notify(PIPE_FIX, PIPE_MTE1, 0)<br>asc_sync_wait(PIPE_FIX, PIPE_MTE1, 0) \| 暂无应用场景 \| 暂无应用场景 \| asc_sync_pipe(PIPE_FIX) \|
	92	+
	93	+<a name="aiv_intra_core_sync_combinations"></a>
	94	+
	95	+表3 AIV中所有合法的核内同步组合
	96	+
	97	+\| 源流水/目的流水 \| PIPE_S \| PIPE_V \| PIPE_MTE2 \| PIPE_MTE3 \|
	98	+\| --- \| --- \| --- \| --- \| --- \|
	99	+\| PIPE_S \| 不涉及 \| asc_sync_notify(PIPE_S, PIPE_V, 0)<br>asc_sync_wait(PIPE_S, PIPE_V, 0) \| asc_sync_notify(PIPE_S, PIPE_MTE2, 0)<br>asc_sync_wait(PIPE_S, PIPE_MTE2, 0) \| asc_sync_notify(PIPE_S, PIPE_MTE3, 0)<br>asc_sync_wait(PIPE_S, PIPE_MTE3, 0) \|
	100	+\| PIPE_V \| asc_sync_notify(PIPE_V, PIPE_S, 0)<br>asc_sync_wait(PIPE_V, PIPE_S, 0) \| asc_sync_pipe(PIPE_V) \| asc_sync_notify(PIPE_V, PIPE_MTE2, 0)<br>asc_sync_wait(PIPE_V, PIPE_MTE2, 0) \| asc_sync_notify(PIPE_V, PIPE_MTE3, 0)<br>asc_sync_wait(PIPE_V, PIPE_MTE3, 0) \|
	101	+\| PIPE_MTE2 \| asc_sync_notify(PIPE_MTE2, PIPE_S, 0)<br>asc_sync_wait(PIPE_MTE2, PIPE_S, 0) \| asc_sync_notify(PIPE_MTE2, PIPE_V, 0)<br>asc_sync_wait(PIPE_MTE2, PIPE_V, 0) \| asc_sync_pipe(PIPE_MTE2) \| asc_sync_notify(PIPE_MTE2, PIPE_MTE3, 0)<br>asc_sync_wait(PIPE_MTE2, PIPE_MTE3, 0) \|
	102	+\| PIPE_MTE3 \| asc_sync_notify(PIPE_MTE3, PIPE_S, 0)<br>asc_sync_wait(PIPE_MTE3, PIPE_S, 0) \| asc_sync_notify(PIPE_MTE3, PIPE_V, 0)<br>asc_sync_wait(PIPE_MTE3, PIPE_V, 0) \| asc_sync_notify(PIPE_MTE3, PIPE_MTE2, 0)<br>asc_sync_wait(PIPE_MTE3, PIPE_MTE2, 0) \| asc_sync_pipe(PIPE_MTE3) \|
	103	+
	104	+<!-- end id4 -->
	105	+
	106	+## 什么时候需要开发者手动插入同步
	107	+
	108	+C API编程方式下所有的同步均需开发者手动管理。

Mdocs/zh/guide/编程指南/编译与运行/算子编译/AI-Core算子编译基本用法.md+1-1

	@@ -113,7 +113,7 @@ AI Core SIMT的基本编译流程如下：Host代码使用Host编译器编译成
113	\| -fPIC \| 否 \| 告知编译器产生位置无关代码。 \|	113	\| -fPIC \| 否 \| 告知编译器产生位置无关代码。 \|
114	\| -O \| 否 \| 用于指定编译器的优化级别，当前支持-O3，-O2，-O0。 \|	114	\| -O \| 否 \| 用于指定编译器的优化级别，当前支持-O3，-O2，-O0。 \|
115	\| --run-mode=sim \| 否 \| sim模式：链接时用户添加仿真模式对应的实现库，实现代码在仿真模式下运行，可以查看仿真相关日志，方便用户性能调试。 \|	115	\| --run-mode=sim \| 否 \| sim模式：链接时用户添加仿真模式对应的实现库，实现代码在仿真模式下运行，可以查看仿真相关日志，方便用户性能调试。 \|
116	-\| --enable-simt \| 否 \| SIMT编程场景，指定SIMT方式编译。 \|	116	+\| --enable-simt \| 否 \| SIMT编程场景，指定SIMT方式编译。设置编译选项`--enable-simt`时，若包含[SIMD API](../../../../api/SIMD-API/SIMD-API.md)的头文件会导致编译失败。 \|
117	\| --cce-disable-asc-reserved-ubuf \| 否 \| 禁用Ascend C接口使用预留UB空间。开启后，依赖预留UB空间的Ascend C接口在对应芯片架构下不可用，使用时编译报错。使用预留UB空间的API列表参考：[使用预留UB空间的API](../../编程模型/AI-Core-SIMD编程/基于Tensor的CPP编程/静态Tensor编程.md#使用预留ub空间的api范围)。由于同一API使用预留UB空间的情况在不同NPU架构下有差异，开启该编译选项后，需要手动调整API调用方式或替换为不依赖预留UB空间的实现，才能完成兼容性迁移。 \|	117	\| --cce-disable-asc-reserved-ubuf \| 否 \| 禁用Ascend C接口使用预留UB空间。开启后，依赖预留UB空间的Ascend C接口在对应芯片架构下不可用，使用时编译报错。使用预留UB空间的API列表参考：[使用预留UB空间的API](../../编程模型/AI-Core-SIMD编程/基于Tensor的CPP编程/静态Tensor编程.md#使用预留ub空间的api范围)。由于同一API使用预留UB空间的情况在不同NPU架构下有差异，开启该编译选项后，需要手动调整API调用方式或替换为不依赖预留UB空间的实现，才能完成兼容性迁移。 \|
118	\| --cce-disable-vf-stack-reserved-ubuf \| 否 \| 禁用SIMD VF栈预留的UB空间。开启后，编译器不再预留该部分UB空间，该空间可作为普通UB空间使用。针对 [NPU架构版本2201](../../语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)，此编译选项无实际效果；针对 [NPU架构版本3510](../../语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)，此编译选项生效，当用户使用此编译选项后，编译器将无法使用预留的UB空间进行寄存器溢出的缓存，需要用户保证寄存器不溢出。 \|	118	\| --cce-disable-vf-stack-reserved-ubuf \| 否 \| 禁用SIMD VF栈预留的UB空间。开启后，编译器不再预留该部分UB空间，该空间可作为普通UB空间使用。针对 [NPU架构版本2201](../../语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)，此编译选项无实际效果；针对 [NPU架构版本3510](../../语言扩展层/SIMD-BuiltIn关键字.md#npu-arch)，此编译选项生效，当用户使用此编译选项后，编译器将无法使用预留的UB空间进行寄存器溢出的缓存，需要用户保证寄存器不溢出。 \|
119	\| --cce-auto-sync \| 否 \| 开启毕昇编译器自动同步。AI Core内部的执行单元是异步并行的，Tensor的读写可能存在数据依赖，开启后可由毕昇编译器自动插入部分同步。详细内容请参考[关键特性说明](../../../../api/SIMD-API/basic_api/sync_control/intra_core_sync/key_features.md)。 \|	119	\| --cce-auto-sync \| 否 \| 开启毕昇编译器自动同步。AI Core内部的执行单元是异步并行的，Tensor的读写可能存在数据依赖，开启后可由毕昇编译器自动插入部分同步。详细内容请参考[关键特性说明](../../../../api/SIMD-API/basic_api/sync_control/intra_core_sync/key_features.md)。 \|