Star784
1081
代码
代码
Issues75
Pull Requests200
Actions
项目讨论
Wiki
分析
Star784
1081
已合并
WriteGmByPassDCache和ReadGmByPassDCache接口名称问题修改 #5327
chenmyk创建于 22 天前
WriteGmByPassDCache和ReadGmByPassDCache接口名称问题修改 #5327
已合并
chenmyk创建于 22 天前
从chenmyk/asc-devkit_7346master_0825_Bypass合入到cann/asc-devkitmaster
共 37 个文件变更+499-193
Mdocs/zh/api/README.md+5-2
@@ -363,6 +363,9 @@
363 - [数据重排（ISASI）](SIMD-API/basic_api/memory_vector_compute/data_rearrange_ISASI/data_rearrange_ISASI.md)363 - [数据重排（ISASI）](SIMD-API/basic_api/memory_vector_compute/data_rearrange_ISASI/data_rearrange_ISASI.md)
364 - [Interleave](SIMD-API/basic_api/memory_vector_compute/data_rearrange_ISASI/Interleave.md)364 - [Interleave](SIMD-API/basic_api/memory_vector_compute/data_rearrange_ISASI/Interleave.md)
365 - [DeInterleave](SIMD-API/basic_api/memory_vector_compute/data_rearrange_ISASI/DeInterleave.md)365 - [DeInterleave](SIMD-API/basic_api/memory_vector_compute/data_rearrange_ISASI/DeInterleave.md)
366+ - [废弃接口](SIMD-API/basic_api/deprecated_interface/deprecated_interface.md)
367+ - [WriteGmByPassDCache\(ISASI\)（废弃）](SIMD-API/basic_api/deprecated_interface/WriteGmByPassDCache_ISASI_deprecated.md)
368+ - [ReadGmByPassDCache\(ISASI\)（废弃）](SIMD-API/basic_api/deprecated_interface/ReadGmByPassDCache_ISASI_deprecated.md)
366 369 
367 - [Reg矢量计算](SIMD-API/basic_api/reg_vector_compute/reg_vector_compute.md)370 - [Reg矢量计算](SIMD-API/basic_api/reg_vector_compute/reg_vector_compute.md)
368 - [概述](SIMD-API/basic_api/reg_vector_compute/overview.md)371 - [概述](SIMD-API/basic_api/reg_vector_compute/overview.md)
@@ -510,8 +513,8 @@
510 - [Cast（float转bfloat16\_t）](SIMD-API/basic_api/scalar_compute/Cast_float_to_bfloat16.md)513 - [Cast（float转bfloat16\_t）](SIMD-API/basic_api/scalar_compute/Cast_float_to_bfloat16.md)
511 - [Cast（多类型转float）](SIMD-API/basic_api/scalar_compute/Cast_multi_type_to_float.md)514 - [Cast（多类型转float）](SIMD-API/basic_api/scalar_compute/Cast_multi_type_to_float.md)
512 - [Nop](SIMD-API/basic_api/scalar_compute/Nop.md)515 - [Nop](SIMD-API/basic_api/scalar_compute/Nop.md)
513- - [WriteGmByPassDCache\(ISASI\)](SIMD-API/basic_api/scalar_compute/WriteGmByPassDCache_ISASI.md)516+ - [WriteGmBypassDCache\(ISASI\)](SIMD-API/basic_api/scalar_compute/WriteGmBypassDCache_ISASI.md)
514- - [ReadGmByPassDCache\(ISASI\)](SIMD-API/basic_api/scalar_compute/ReadGmByPassDCache_ISASI.md)517+ - [ReadGmBypassDCache\(ISASI\)](SIMD-API/basic_api/scalar_compute/ReadGmBypassDCache_ISASI.md)
515 518 
516 - [资源管理](SIMD-API/basic_api/resource_management/resource_management.md)519 - [资源管理](SIMD-API/basic_api/resource_management/resource_management.md)
517 - [TPipe](SIMD-API/basic_api/resource_management/TPipe/TPipe.md)520 - [TPipe](SIMD-API/basic_api/resource_management/TPipe/TPipe.md)
Mdocs/zh/api/SIMD-API/basic_api/atomic_operations/key_features.md+27-23
@@ -16,14 +16,16 @@
16 16 
17如下图1所示，开发者的预期结果：指令发射的顺序能够严格对应实际指令执行顺序，多次执行该段代码，无论执行多少次，最终GM数据均为data0 + data1 + data2 + data3，结果完全一致，实现确定性计算。17如下图1所示，开发者的预期结果：指令发射的顺序能够严格对应实际指令执行顺序，多次执行该段代码，无论执行多少次，最终GM数据均为data0 + data1 + data2 + data3，结果完全一致，实现确定性计算。
18 18 
19-**图1** 确定性计算场景，GM上数据变化过程<a name="zh-cn_topic_0000002552941426_fig71821376818"></a> 19+**图1** 确定性计算场景，GM上数据变化过程<a name="zh-cn_topic_0000002552941426_fig71821376818"></a>
20+ 
20![](../../../figures/deterministic_computation_gm_data_change_process.png "Deterministic computation scenario, GM data change process")21![](../../../figures/deterministic_computation_gm_data_change_process.png "Deterministic computation scenario, GM data change process")
21 22 
22但实际情况是，若开发者不做干预，程序每次运行时这些指令的执行顺序都可能发生变化，最终导致GM数据与预期结果不一致。下面列举两种可能的指令执行顺序及其对应的执行流程。23但实际情况是，若开发者不做干预，程序每次运行时这些指令的执行顺序都可能发生变化，最终导致GM数据与预期结果不一致。下面列举两种可能的指令执行顺序及其对应的执行流程。
23 24 
24## 非确定性计算，结果1<a name="zh-cn_topic_0000002552941426_section027603463315"></a>25## 非确定性计算，结果1<a name="zh-cn_topic_0000002552941426_section027603463315"></a>
25 26 
26-**图2** 非确定性计算场景1，GM上数据变化过程<a name="zh-cn_topic_0000002552941426_fig273814664819"></a> 27+**图2** 非确定性计算场景1，GM上数据变化过程<a name="zh-cn_topic_0000002552941426_fig273814664819"></a>
28+ 
27![](../../../figures/non_deterministic_computation_1_gm_data_change_process.png "Non-deterministic computation scenario 1, GM data change process")29![](../../../figures/non_deterministic_computation_1_gm_data_change_process.png "Non-deterministic computation scenario 1, GM data change process")
28 30 
29如图2所示，该场景中指令执行流程如下：31如图2所示，该场景中指令执行流程如下：
@@ -45,7 +47,8 @@
45 47 
46## 非确定性计算，结果2<a name="zh-cn_topic_0000002552941426_section1547411390818"></a>48## 非确定性计算，结果2<a name="zh-cn_topic_0000002552941426_section1547411390818"></a>
47 49 
48-**图3** 非确定性计算场景2，GM上数据变化过程<a name="zh-cn_topic_0000002552941426_fig1829915169912"></a> 50+**图3** 非确定性计算场景2，GM上数据变化过程<a name="zh-cn_topic_0000002552941426_fig1829915169912"></a>
51+ 
49![](../../../figures/non_deterministic_computation_2_gm_data_change_process.png "Non-deterministic computation scenario 2, GM data change process")52![](../../../figures/non_deterministic_computation_2_gm_data_change_process.png "Non-deterministic computation scenario 2, GM data change process")
50 53 
51如图3所示，该场景中指令执行流程如下：54如图3所示，该场景中指令执行流程如下：
@@ -98,20 +101,20 @@ if (GetBlockIdx() == 0) {
98 向GM搬运数据data0;101 向GM搬运数据data0;
99 核间同步102 核间同步
100} else if (GetBlockIdx() == 1) {103} else if (GetBlockIdx() == 1) {
101- 核间同步 104+ 核间同步
102- SetAtomicAdd(); 105+ SetAtomicAdd();
103 向GM搬运data1;106 向GM搬运data1;
104- 核间同步 107+ 核间同步
105} else if (GetBlockIdx() == 2) {108} else if (GetBlockIdx() == 2) {
106 核间同步109 核间同步
107- SetAtomicAdd(); 110+ SetAtomicAdd();
108- 向GM搬运data2; 111+ 向GM搬运data2;
109- 核间同步 112+ 核间同步
110} else if (GetBlockIdx() == 3) {113} else if (GetBlockIdx() == 3) {
111 核间同步114 核间同步
112- SetAtomicAdd(); 115+ SetAtomicAdd();
113- 向GM搬运data3; 116+ 向GM搬运data3;
114-} 117+}
115```118```
116 119 
117下面介绍如何基于硬件同步指令实现核内同步，以及如何基于软件同步方案实现核间同步。120下面介绍如何基于硬件同步指令实现核内同步，以及如何基于软件同步方案实现核间同步。
@@ -149,7 +152,8 @@ if (GetBlockIdx() == 0) {
149 152 
150- 当前核在执行搬运任务前，会通过Scalar单元不断读取该信号量的值。如果信号量不等于1，当前核会进入阻塞等待状态；当检测到信号量等于1时，当前核会解除阻塞，开始执行自己的数据搬运或原子操作。为确保信号量等于1之前，当前核不会执行搬运指令，需要在搬运指令之前插入核内同步3。153- 当前核在执行搬运任务前，会通过Scalar单元不断读取该信号量的值。如果信号量不等于1，当前核会进入阻塞等待状态；当检测到信号量等于1时，当前核会解除阻塞，开始执行自己的数据搬运或原子操作。为确保信号量等于1之前，当前核不会执行搬运指令，需要在搬运指令之前插入核内同步3。
151 154 
152-**图4** 一对核之间软件同步方案流程图<a name="zh-cn_topic_0000002583421469_fig1491417244811"></a> 155+**图4** 一对核之间软件同步方案流程图<a name="zh-cn_topic_0000002583421469_fig1491417244811"></a>
156+ 
153![](../../../figures/software_synchronization_between_two_cores_flowchart.png "Software synchronization flowchart between two cores")157![](../../../figures/software_synchronization_between_two_cores_flowchart.png "Software synchronization flowchart between two cores")
154 158 
155Scalar单元访问GM上的信号量，存在两种访问方式：159Scalar单元访问GM上的信号量，存在两种访问方式：
@@ -160,7 +164,7 @@ Scalar单元访问GM上的信号量，存在两种访问方式：
160 164 
161- 不通过DCache访问165- 不通过DCache访问
162 166 
163- 使用[WriteGmByPassDCache](../scalar_compute/WriteGmByPassDCache_ISASI.md)和[ReadGmByPassDCache](../scalar_compute/ReadGmByPassDCache_ISASI.md)。这种方式无需额外操作即可保证多核间数据的一致性。167+ 使用[WriteGmBypassDCache](../scalar_compute/WriteGmBypassDCache_ISASI.md)和[ReadGmBypassDCache](../scalar_compute/ReadGmBypassDCache_ISASI.md)。这种方式无需额外操作即可保证多核间数据的一致性。
164 168 
165两种方案在性能上的差异：不经过DCache，性能会较差，但是如果读写GM数据较少，可以考虑使用不经过DCache的方法。169两种方案在性能上的差异：不经过DCache，性能会较差，但是如果读写GM数据较少，可以考虑使用不经过DCache的方法。
166 170 
@@ -294,15 +298,15 @@ Scalar单元访问GM上的信号量，存在两种访问方式：
294 // 在通知下一个核之前，等待当前核的任务完成。298 // 在通知下一个核之前，等待当前核的任务完成。
295 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);299 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);
296 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);300 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);
297- AscendC::WriteGmByPassDCache<int32_t>(addr + blockIdx * 32, 1);301+ AscendC::WriteGmBypassDCache<int32_t>(addr + blockIdx * 32, 1);
298 } else if (blockIdx == 1) {302 } else if (blockIdx == 1) {
299 int32_t preblockIdx = 3;303 int32_t preblockIdx = 3;
300 304 
301 // 核1等待核3，直到核3数据已完成累加。305 // 核1等待核3，直到核3数据已完成累加。
302 while (true) {306 while (true) {
303- int32_t value = AscendC::ReadGmByPassDCache<int32_t>(addr + preblockIdx * 32);307+ int32_t value = AscendC::ReadGmBypassDCache<int32_t>(addr + preblockIdx * 32);
304 if (value == 1) {308 if (value == 1) {
305- AscendC::WriteGmByPassDCache<int32_t>(addr + preblockIdx * 32, 0);309+ AscendC::WriteGmBypassDCache<int32_t>(addr + preblockIdx * 32, 0);
306 break;310 break;
307 }311 }
308 }312 }
@@ -323,9 +327,9 @@ Scalar单元访问GM上的信号量，存在两种访问方式：
323 327 
324 // 核2等待核0，直到核0已完成GM初始化。328 // 核2等待核0，直到核0已完成GM初始化。
325 while (true) {329 while (true) {
326- int32_t value = AscendC::ReadGmByPassDCache<int32_t>(addr + preblockIdx * 32);330+ int32_t value = AscendC::ReadGmBypassDCache<int32_t>(addr + preblockIdx * 32);
327 if (value == 1) {331 if (value == 1) {
328- AscendC::WriteGmByPassDCache<int32_t>(addr + preblockIdx * 32, 0);332+ AscendC::WriteGmBypassDCache<int32_t>(addr + preblockIdx * 32, 0);
329 break;333 break;
330 }334 }
331 }335 }
@@ -345,15 +349,15 @@ Scalar单元访问GM上的信号量，存在两种访问方式：
345 // 在通知下一个核之前，等待当前核的任务完成。349 // 在通知下一个核之前，等待当前核的任务完成。
346 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);350 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);
347 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);351 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);
348- AscendC::WriteGmByPassDCache<int32_t>(addr + blockIdx * 32, 1);352+ AscendC::WriteGmBypassDCache<int32_t>(addr + blockIdx * 32, 1);
349 } else if (blockIdx == 3) {353 } else if (blockIdx == 3) {
350 int32_t preblockIdx = 2;354 int32_t preblockIdx = 2;
351 355 
352 // 核3等待核2，直到核2数据已完成累加。356 // 核3等待核2，直到核2数据已完成累加。
353 while (true) {357 while (true) {
354- int32_t value = AscendC::ReadGmByPassDCache<int32_t>(addr + preblockIdx * 32);358+ int32_t value = AscendC::ReadGmBypassDCache<int32_t>(addr + preblockIdx * 32);
355 if (value == 1) {359 if (value == 1) {
356- AscendC::WriteGmByPassDCache<int32_t>(addr + preblockIdx * 32, 0);360+ AscendC::WriteGmBypassDCache<int32_t>(addr + preblockIdx * 32, 0);
357 break;361 break;
358 }362 }
359 }363 }
@@ -374,7 +378,7 @@ Scalar单元访问GM上的信号量，存在两种访问方式：
374 // 在通知下一个核之前，等待当前核的任务完成。378 // 在通知下一个核之前，等待当前核的任务完成。
375 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);379 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);
376 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);380 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);
377- AscendC::WriteGmByPassDCache<int32_t>(addr + blockIdx * 32, 1);381+ AscendC::WriteGmBypassDCache<int32_t>(addr + blockIdx * 32, 1);
378 }382 }
379 // GM结果：data0 + data2 + data3 + data1。383 // GM结果：data0 + data2 + data3 + data1。
380 ```384 ```
Mdocs/zh/api/SIMD-API/basic_api/basic_api_list.md+2-2
@@ -448,8 +448,8 @@
448| [Cast（float转bfloat16_t）](scalar_compute/Cast_float_to_bfloat16.md) | float类型标量数据转换成bfloat16_t类型标量数据。 |448| [Cast（float转bfloat16_t）](scalar_compute/Cast_float_to_bfloat16.md) | float类型标量数据转换成bfloat16_t类型标量数据。 |
449| [Cast（多类型转float）](scalar_compute/Cast_multi_type_to_float.md) | 该接口将输入数据转换为float类型。 |449| [Cast（多类型转float）](scalar_compute/Cast_multi_type_to_float.md) | 该接口将输入数据转换为float类型。 |
450| [Nop](scalar_compute/Nop.md) | 用户输入延迟的个数，实现延迟若干个cycle。 |450| [Nop](scalar_compute/Nop.md) | 用户输入延迟的个数，实现延迟若干个cycle。 |
451-| [WriteGmByPassDCache(ISASI)](scalar_compute/WriteGmByPassDCache_ISASI.md) | 不经过DCache向GM地址上写数据。 |451+| [WriteGmBypassDCache(ISASI)](scalar_compute/WriteGmBypassDCache_ISASI.md) | 不经过DCache向GM地址上写数据。 |
452-| [ReadGmByPassDCache(ISASI)](scalar_compute/ReadGmByPassDCache_ISASI.md) | 不经过DCache从GM地址上读数据。 |452+| [ReadGmBypassDCache(ISASI)](scalar_compute/ReadGmBypassDCache_ISASI.md) | 不经过DCache从GM地址上读数据。 |
453 453 
454## 资源管理454## 资源管理
455| 接口名 | 功能描述 |455| 接口名 | 功能描述 |
Mdocs/zh/api/SIMD-API/basic_api/cache_control/system_cache_overview.md+6-3
@@ -8,13 +8,16 @@ Cache（缓存）的主要作用是在搬运单元或Scalar单元与外部存储
8<!-- npu="950" id1 -->8<!-- npu="950" id1 -->
9以[NPU架构版本3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)为例，图1展示了AI Core中支持的五类Cache（L2 Cache、DCache、ICache、SIMT DCache、NDDMA Cache）在硬件架构中的位置关系。9以[NPU架构版本3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)为例，图1展示了AI Core中支持的五类Cache（L2 Cache、DCache、ICache、SIMT DCache、NDDMA Cache）在硬件架构中的位置关系。
10 10 
11-**图1** 五类Cache在AI Core中的位置关系示意图 11+**图1** 五类Cache在AI Core中的位置关系示意图
M
Mmunanhw15 天前

不要删除空格，排查全文。 删除图标题和图之间的空行或者空格 回导致图片和标题显示在一行

likedislike
12+ 
12![](../../../../api/figures/npu_3510_hw_arch_cache.png "atlas_950_cache_architecture_diagram")13![](../../../../api/figures/npu_3510_hw_arch_cache.png "atlas_950_cache_architecture_diagram")
14+ 
13<!-- end id1 -->15<!-- end id1 -->
14<!-- npu="A3,910b" id2 -->16<!-- npu="A3,910b" id2 -->
15以[NPU架构版本2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)为例，图2展示了AI Core中支持的三类Cache（L2 Cache、DCache和ICache）在硬件架构中的位置关系。17以[NPU架构版本2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)为例，图2展示了AI Core中支持的三类Cache（L2 Cache、DCache和ICache）在硬件架构中的位置关系。
16 18 
17-**图2** 三类Cache在AI Core中的位置关系示意图<a name="zh-cn_topic_0000002583420201_fig496311278412"></a> 19+**图2** 三类Cache在AI Core中的位置关系示意图<a name="zh-cn_topic_0000002583420201_fig496311278412"></a>
20+ 
18![](../../../../api/figures/atlas_a2_a3_cache_architecture_diagram.png "atlas_a2_a3_cache_architecture_diagram")21![](../../../../api/figures/atlas_a2_a3_cache_architecture_diagram.png "atlas_a2_a3_cache_architecture_diagram")
19<!-- end id2 -->22<!-- end id2 -->
20 23 
@@ -97,7 +100,7 @@ Cache一致性是多核中确保数据正确性的核心机制。简单来说，
97| L2 Cache | [SetL2CacheHint](../../basic_api/data_structures/GlobalTensor/SetL2CacheHint.md) | 读写GM的数据都默认被缓存在L2 Cache（默认启用L2 Cache），SetL2CacheHint能够设置GlobalTensor是否启用L2 Cache。 |100| L2 Cache | [SetL2CacheHint](../../basic_api/data_structures/GlobalTensor/SetL2CacheHint.md) | 读写GM的数据都默认被缓存在L2 Cache（默认启用L2 Cache），SetL2CacheHint能够设置GlobalTensor是否启用L2 Cache。 |
98| DCache | [DataCachePreload](DataCachePreload.md) | DataCachePreload接口从源地址所在的特定GM地址预加载数据到DCache中，每次调用只能预加载一个Cache Line大小的数据。 |101| DCache | [DataCachePreload](DataCachePreload.md) | DataCachePreload接口从源地址所在的特定GM地址预加载数据到DCache中，每次调用只能预加载一个Cache Line大小的数据。 |
99| DCache | [DataCacheCleanAndInvalid](DataCacheCleanAndInvalid.md) | 当Scalar单元访问GM时，使用该接口刷新Cache，保证Cache的一致性，使用场景如下：<br>&bull;读取GM的数据，但该数据可能在外部被其余核修改，此时需要使用DataCacheCleanAndInvalid接口，直接访问GM，获取最新数据。<br>&bull;用户通过Scalar单元写GM的数据，希望立刻写出，也需要使用DataCacheCleanAndInvalid接口。 |102| DCache | [DataCacheCleanAndInvalid](DataCacheCleanAndInvalid.md) | 当Scalar单元访问GM时，使用该接口刷新Cache，保证Cache的一致性，使用场景如下：<br>&bull;读取GM的数据，但该数据可能在外部被其余核修改，此时需要使用DataCacheCleanAndInvalid接口，直接访问GM，获取最新数据。<br>&bull;用户通过Scalar单元写GM的数据，希望立刻写出，也需要使用DataCacheCleanAndInvalid接口。 |
100-| DCache | [ReadGmByPassDCache](../scalar_compute/ReadGmByPassDCache_ISASI.md)/[WriteGmByPassDCache](../scalar_compute/WriteGmByPassDCache_ISASI.md) | 不经过DCache向GM地址上**读**/**写**数据。使用场景：<br>&bull;当多个核写入的数据落在同一条Cache Line内时，经过DCache的读写将以64B为粒度，可能引发多核数据随机覆盖问题。使用该接口不经过DCache直接按操作数大小**读**/**写**GM，可避免此问题。<br>&bull;使用该接口不经过DCache直接向GM**读**/**写**数据，可避免由DCache缓存引发的多核间数据不一致问题。 |103+| DCache | [ReadGmBypassDCache](../scalar_compute/ReadGmBypassDCache_ISASI.md)/[WriteGmBypassDCache](../scalar_compute/WriteGmBypassDCache_ISASI.md) | 不经过DCache向GM地址上**读**/**写**数据。使用场景：<br>&bull;当多个核写入的数据落在同一条Cache Line内时，经过DCache的读写将以64B为粒度，可能引发多核数据随机覆盖问题。使用该接口不经过DCache直接按操作数大小**读**/**写**GM，可避免此问题。<br>&bull;使用该接口不经过DCache直接向GM**读**/**写**数据，可避免由DCache缓存引发的多核间数据不一致问题。 |
101| ICache | [ICachePreLoad](ICachePreLoad_ISASI.md) | 开发者手动调用ICachePreLoad接口，能够从指令所在GM地址预加载指令到ICache中。 |104| ICache | [ICachePreLoad](ICachePreLoad_ISASI.md) | 开发者手动调用ICachePreLoad接口，能够从指令所在GM地址预加载指令到ICache中。 |
102| ICache | [GetICachePreloadStatus](GetICachePreloadStatus_ISASI.md) | GetICachePreloadStatus为调试接口，在调用ICachePreLoad后调用，用于获取ICache的PreLoad的状态。当返回值为0时，说明ICache的PreLoad已完成；当返回值为1时，说明ICache的PreLoad未完成。 |105| ICache | [GetICachePreloadStatus](GetICachePreloadStatus_ISASI.md) | GetICachePreloadStatus为调试接口，在调用ICachePreLoad后调用，用于获取ICache的PreLoad的状态。当返回值为0时，说明ICache的PreLoad已完成；当返回值为1时，说明ICache的PreLoad未完成。 |
103| NDDMA Cache | NdDmaDci | 在使用[DataCopy（GMToUB多维数据搬运NDDMA）](../../basic_api/memory_vector_compute/data_move/DataCopy_GMToUB_NDDMA.md)接口进行数据搬运前，需要使用NdDmaDci接口刷新NDDMA Cache缓存保证为最新状态。 |106| NDDMA Cache | NdDmaDci | 在使用[DataCopy（GMToUB多维数据搬运NDDMA）](../../basic_api/memory_vector_compute/data_move/DataCopy_GMToUB_NDDMA.md)接口进行数据搬运前，需要使用NdDmaDci接口刷新NDDMA Cache缓存保证为最新状态。 |
Rdocs/zh/api/SIMD-API/basic_api/scalar_compute/ReadGmByPassDCache_ISASI.md→docs/zh/api/SIMD-API/basic_api/deprecated_interface/ReadGmByPassDCache_ISASI_deprecated.md+9-7
@@ -1,4 +1,4 @@
1-# ReadGmByPassDCache\(ISASI\)<a name="ZH-CN_TOPIC_0000002327311338"></a>1+# ReadGmByPassDCache\(ISASI\)（废弃）
2 2 
3## 产品支持情况3## 产品支持情况
4 4 
@@ -25,20 +25,22 @@
25<!-- end id7 -->25<!-- end id7 -->
26<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/scalar_compute/ReadGmByPassDCache_ISASI_res.md#id1 -->26<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/scalar_compute/ReadGmByPassDCache_ISASI_res.md#id1 -->
27 27 
28-## 功能说明<a name="section618mcpsimp"></a>28+## 功能说明
29+ 
30+**该接口已废弃，请使用[ReadGmBypassDCache\(ISASI\)](../scalar_compute/ReadGmBypassDCache_ISASI.md)接口替代。**
29 31 
30头文件路径为：`"basic_api/kernel_operator_scalar_intf.h"`。32头文件路径为：`"basic_api/kernel_operator_scalar_intf.h"`。
31 33 
32不经过DCache从GM地址上读数据。34不经过DCache从GM地址上读数据。
33 35 
34-## 函数原型<a name="section620mcpsimp"></a>36+## 函数原型
35 37 
36```cpp38```cpp
37template <typename T>39template <typename T>
38__aicore__ inline T ReadGmByPassDCache(__gm__ T* addr)40__aicore__ inline T ReadGmByPassDCache(__gm__ T* addr)
39```41```
40 42 
41-## 参数说明<a name="section622mcpsimp"></a>43+## 参数说明
42 44 
43**表1** 模板参数说明45**表1** 模板参数说明
44 46 
@@ -56,15 +58,15 @@ __aicore__ inline T ReadGmByPassDCache(__gm__ T* addr)
56 58 
57支持的数据类型为int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。59支持的数据类型为int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。
58 60 
59-## 返回值说明<a name="section640mcpsimp"></a>61+## 返回值说明
60 62 
61源GM地址上的数据。63源GM地址上的数据。
62 64 
63-## 约束说明<a name="section633mcpsimp"></a>65+## 约束说明
64 66 
65无67无
66 68 
67-## 调用示例<a name="section6191129670"></a>69+## 调用示例
68 70 
69```cpp71```cpp
70if (blockIdx == 0) {72if (blockIdx == 0) {
Rdocs/zh/api/SIMD-API/basic_api/scalar_compute/WriteGmByPassDCache_ISASI.md→docs/zh/api/SIMD-API/basic_api/deprecated_interface/WriteGmByPassDCache_ISASI_deprecated.md+10-8
@@ -1,4 +1,4 @@
1-# WriteGmByPassDCache\(ISASI\)<a name="ZH-CN_TOPIC_0000002327151478"></a>1+# WriteGmByPassDCache\(ISASI\)（废弃）
2 2 
3## 产品支持情况3## 产品支持情况
4 4 
@@ -25,13 +25,15 @@
25<!-- end id7 -->25<!-- end id7 -->
26<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/scalar_compute/WriteGmByPassDCache_ISASI_res.md#id1 -->26<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/scalar_compute/WriteGmByPassDCache_ISASI_res.md#id1 -->
27 27 
28-## 功能说明<a name="section618mcpsimp"></a>28+## 功能说明
29+ 
30+**该接口已废弃，请使用[WriteGmBypassDCache\(ISASI\)](../scalar_compute/WriteGmBypassDCache_ISASI.md)接口替代。**
29 31 
30头文件路径为：`"basic_api/kernel_operator_scalar_intf.h"`。32头文件路径为：`"basic_api/kernel_operator_scalar_intf.h"`。
31 33 
32不经过DCache向GM地址上写数据。34不经过DCache向GM地址上写数据。
33<!-- npu="A3,910b" id8 -->35<!-- npu="A3,910b" id8 -->
34-> [!CAUTION]注意 36+> [!CAUTION]注意
35> 针对[NPU架构2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)，接口能否将value成功写入GM还与目标地址addr有关，具体请参见[约束说明](#约束说明)。37> 针对[NPU架构2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)，接口能否将value成功写入GM还与目标地址addr有关，具体请参见[约束说明](#约束说明)。
36<!-- end id8 -->38<!-- end id8 -->
37 39 
@@ -54,14 +56,14 @@ AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
54AscendC::WriteGmByPassDCache<T>(addr, value);56AscendC::WriteGmByPassDCache<T>(addr, value);
55```57```
56 58 
57-## 函数原型<a name="section620mcpsimp"></a>59+## 函数原型
58 60 
59```cpp61```cpp
60template <typename T>62template <typename T>
61__aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value)63__aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value)
62```64```
63 65 
64-## 参数说明<a name="section622mcpsimp"></a>66+## 参数说明
65 67 
66**表1** 模板参数说明68**表1** 模板参数说明
67 69 
@@ -80,11 +82,11 @@ __aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value)
80 82 
81支持的数据类型为int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。83支持的数据类型为int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。
82 84 
83-## 返回值说明<a name="section640mcpsimp"></a>85+## 返回值说明
84 86 
85无87无
86 88 
87-## 约束说明<a name="section633mcpsimp"></a>89+## 约束说明
88 90 
89<!-- npu="A3,910b" id9 -->91<!-- npu="A3,910b" id9 -->
90针对[NPU架构2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)，接口是否执行写入取决于目标地址addr在当前128字节对齐区间内的偏移。令$\mathrm{offset} = addr \bmod 128$，当$0 \leq \mathrm{offset} < 32$时，写入生效；当$32 \leq \mathrm{offset} < 128$时，不执行写入，目标地址中的数据保持原值。即仅当目标地址位于每个128字节对齐区间的前32字节时，接口才执行写入。92针对[NPU架构2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)，接口是否执行写入取决于目标地址addr在当前128字节对齐区间内的偏移。令$\mathrm{offset} = addr \bmod 128$，当$0 \leq \mathrm{offset} < 32$时，写入生效；当$32 \leq \mathrm{offset} < 128$时，不执行写入，目标地址中的数据保持原值。即仅当目标地址位于每个128字节对齐区间的前32字节时，接口才执行写入。
@@ -100,7 +102,7 @@ __aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value)
100 $$102 $$
101<!-- end id9 -->103<!-- end id9 -->
102 104 
103-## 调用示例<a name="section6191129670"></a>105+## 调用示例
104 106 
105```cpp107```cpp
106if (blockIdx == 0) {108if (blockIdx == 0) {
Adocs/zh/api/SIMD-API/basic_api/deprecated_interface/deprecated_interface.md+7-0
@@ -0,0 +1,7 @@
1+# 废弃接口
2+ 
3+## 标量计算
4+ 
5+- **[WriteGmByPassDCache\(ISASI\)（废弃）](WriteGmByPassDCache_ISASI_deprecated.md)**
6+ 
7+- **[ReadGmByPassDCache\(ISASI\)（废弃）](ReadGmByPassDCache_ISASI_deprecated.md)**
Adocs/zh/api/SIMD-API/basic_api/scalar_compute/ReadGmBypassDCache_ISASI.md+93-0
@@ -0,0 +1,93 @@
1+# ReadGmBypassDCache\(ISASI\)
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- Ascend 950PR/Ascend 950DT：支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- Atlas A3 训练系列产品/Atlas A3 推理系列产品：支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- Atlas A2 训练系列产品/Atlas A2 推理系列产品：支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- Atlas 200I/500 A2 推理产品：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- Atlas 推理系列产品AI Core：不支持
19+<!-- end id5 -->
20+<!-- npu="310p" id6 -->
21+- Atlas 推理系列产品Vector Core：不支持
22+<!-- end id6 -->
23+<!-- npu="910" id7 -->
24+- Atlas 训练系列产品：不支持
25+<!-- end id7 -->
26+ 
27+## 功能说明
28+ 
29+头文件路径为：`"basic_api/kernel_operator_scalar_intf.h"`。
30+ 
31+不经过DCache从GM地址上读数据。
32+ 
33+## 函数原型
34+ 
35+```cpp
36+template <typename T>
37+__aicore__ inline T ReadGmBypassDCache(__gm__ T* addr)
38+```
39+ 
40+## 参数说明
41+ 
42+**表1** 模板参数说明
43+ 
44+| 参数名 | 描述 |
45+|--------|------|
46+| T | 操作数的数据类型。 |
47+ 
48+**表2** 接口参数说明
49+ 
50+| 参数名 | 输入/输出 | 含义 |
51+|--------|-----------|------|
52+| addr | 输入 | 源GM地址。 |
53+ 
54+## 数据类型
55+ 
56+支持的数据类型为int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。
57+ 
58+## 返回值说明
59+ 
60+源GM地址上的数据。
61+ 
62+## 约束说明
63+ 
64+无
65+ 
66+## 调用示例
67+ 
68+```cpp
69+if (blockIdx == 0) {
70+ // 先写入被依赖数据，再通过DataSyncBarrier等待DDR访问完成。
71+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm) + 1, DATA_VALUE);
72+ // DataSyncBarrier<DDR>阻塞后续GM写，确保上一条GM写对其他核可见。
73+ AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
74+ // 最后写入同步标记，block 1读到该标记后即可安全读取srcGm[1]。
75+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm), SYNC_FLAG);
76+}
77+ 
78+if (blockIdx == 1) {
79+ while (true) {
80+ __gm__ T *addr = const_cast<__gm__ T *>(srcGlobal.GetPhyAddr());
81+ // 轮询GM第0个元素，等待block 0写入同步标记。
82+ T flagValue = AscendC::ReadGmBypassDCache<T>(addr);
83+ if (flagValue == SYNC_FLAG) {
84+ // DataSyncBarrier保证同步标记之前的srcGm[1]写入已完成。
85+ T dataValue = AscendC::ReadGmBypassDCache<T>(addr + 1);
86+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(dstGm), 2 * dataValue);
87+ return;
88+ }
89+ }
90+}
91+```
92+ 
93+完整样例请参考[DataSyncBarrier样例](../../../../../../examples/01_simd_cpp_api/03_basic_api/05_sync_control/data_sync_barrier)。
Adocs/zh/api/SIMD-API/basic_api/scalar_compute/WriteGmBypassDCache_ISASI.md+129-0
@@ -0,0 +1,129 @@
1+# WriteGmBypassDCache\(ISASI\)
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- Ascend 950PR/Ascend 950DT：支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- Atlas A3 训练系列产品/Atlas A3 推理系列产品：支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- Atlas A2 训练系列产品/Atlas A2 推理系列产品：支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- Atlas 200I/500 A2 推理产品：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- Atlas 推理系列产品AI Core：不支持
19+<!-- end id5 -->
20+<!-- npu="310p" id6 -->
21+- Atlas 推理系列产品Vector Core：不支持
22+<!-- end id6 -->
23+<!-- npu="910" id7 -->
24+- Atlas 训练系列产品：不支持
25+<!-- end id7 -->
26+ 
27+## 功能说明
28+ 
29+头文件路径为：`"basic_api/kernel_operator_scalar_intf.h"`。
30+ 
31+不经过DCache向GM地址上写数据。
32+<!-- npu="A3,910b" id8 -->
33+> [!CAUTION]注意
34+> 针对[NPU架构2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)，接口能否将value成功写入GM还与目标地址addr有关，具体请参见[约束说明](#约束说明)。
35+<!-- end id8 -->
36+ 
37+使用场景：
38+ 
39+- 当多个核写入的数据落在同一条Cache Line内时，经过DCache的读写将以64B为粒度，可能引发多核数据随机覆盖问题（参考[DataCacheCleanAndInvalid调用示例3](../cache_control/DataCacheCleanAndInvalid.md#example3_multi_core)）。使用该接口不经过DCache直接按操作数大小写GM，可避免此问题。
40+- 经过DCache写GM时可能导致多核间的数据不一致问题（详细原因请参考[Cache写策略与Cache一致性问题](../cache_control/system_cache_overview.md#zh-cn_topic_0000002583420201_section053731716357)），使用该接口不经过DCache直接向GM写数据，可避免此问题。
41+ 
42+由于WriteGmBypassDCache接口向GM写数据时不经过DCache，因此开发者需要考虑如下场景：
43+- 当GM上地址addr已经在DCache缓存并且其对应的Cache Line标记为“脏”（dirty，表示该数据已被修改但尚未写回到GM）时，开发者应该在WriteGmBypassDCache接口之前调用[DataCacheCleanAndInvalid](../cache_control/DataCacheCleanAndInvalid.md)接口将addr对应的Cache Line立即写回GM，否则WriteGmBypassDCache接口写入addr的数据后续可能被DCache写回的脏数据覆盖。
44+- DataCacheCleanAndInvalid接口与WriteGmBypassDCache接口向GM写数据时硬件不能保证两个接口的执行顺序，因此开发者应该在WriteGmBypassDCache接口之前调用[DataSyncBarrier](../sync_control/intra_core_sync/DataSyncBarrier_ISASI.md)接口对这两个接口进行同步，否则WriteGmBypassDCache接口写入addr的数据依然可能被DCache写回的脏数据覆盖。
45+ 
46+根据以上的描述，为了简化编程（开发者无需关心addr是否在DCache缓存以及是否被标记为“脏”），建议开发者在使用WriteGmBypassDCache接口时采用如下代码片段：
47+ 
48+```cpp
49+AscendC::GlobalTensor<T> global; // global为addr对应的GlobalTensor。
50+AscendC::DataCacheCleanAndInvalid<T, AscendC::CacheLine::SINGLE_CACHE_LINE, AscendC::DcciDst::CACHELINE_OUT>(global);
51+// 保证WriteGmBypassDCache接口向addr写入value之前，DCache中的脏数据已经写回GM。
52+AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
53+AscendC::WriteGmBypassDCache<T>(addr, value);
54+```
55+ 
56+## 函数原型
57+ 
58+```cpp
59+template <typename T>
60+__aicore__ inline void WriteGmBypassDCache(__gm__ T* addr, T value)
61+```
62+ 
63+## 参数说明
64+ 
65+**表1** 模板参数说明
66+ 
67+| 参数名 | 描述 |
68+|--------|------|
69+| T | 操作数的数据类型。 |
70+ 
71+**表2** 接口参数说明
72+ 
73+| 参数名 | 输入/输出 | 含义 |
74+|--------|-----------|------|
75+| addr | 输入 | 目标GM地址。 |
76+| value | 输入 | 待写入目标数据。 |
77+ 
78+## 数据类型
79+ 
80+支持的数据类型为int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。
81+ 
82+## 返回值说明
83+ 
84+无
85+ 
86+## 约束说明
87+ 
88+<!-- npu="A3,910b" id9 -->
89+针对[NPU架构2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)，接口是否执行写入取决于目标地址addr在当前128字节对齐区间内的偏移。令$\mathrm{offset} = addr \bmod 128$，当$0 \leq \mathrm{offset} < 32$时，写入生效；当$32 \leq \mathrm{offset} < 128$时，不执行写入，目标地址中的数据保持原值。即仅当目标地址位于每个128字节对齐区间的前32字节时，接口才执行写入。当开启AscendC调试宏[-DASCENDC_DEBUG](../../../../guide/programming_guide/compilation_and_execution/operator_compilation/ai_core_operator_compilation.md#内置编译宏开关)时，会有assert校验生效，校验不通过会有assert日志打屏。
90+ 
91+设接口调用前后的目标地址数据分别为$\mathrm{GM}_{\mathrm{before}}(addr)$和$\mathrm{GM}_{\mathrm{after}}(addr)$，则：
92+ 
93+ $$
94+ \mathrm{GM}_{\mathrm{after}}(addr) =
95+ \begin{cases}
96+ value, & 0 \leq \mathrm{offset} < 32 \\
97+ \mathrm{GM}_{\mathrm{before}}(addr), & 32 \leq \mathrm{offset} < 128
98+ \end{cases}
99+ $$
100+<!-- end id9 -->
101+ 
102+## 调用示例
103+ 
104+```cpp
105+if (blockIdx == 0) {
106+ // 先写入被依赖数据，再通过DataSyncBarrier等待DDR访问完成。
107+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm) + 1, DATA_VALUE);
108+ // DataSyncBarrier<DDR>阻塞后续GM写，确保上一条GM写对其他核可见。
109+ AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
110+ // 最后写入同步标记，block 1读到该标记后即可安全读取srcGm[1]。
111+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm), SYNC_FLAG);
112+}
113+ 
114+if (blockIdx == 1) {
115+ while (true) {
116+ __gm__ T *addr = const_cast<__gm__ T *>(srcGlobal.GetPhyAddr());
117+ // 轮询GM第0个元素，等待block 0写入同步标记。
118+ T flagValue = AscendC::ReadGmBypassDCache<T>(addr);
119+ if (flagValue == SYNC_FLAG) {
120+ // DataSyncBarrier保证同步标记之前的srcGm[1]写入已完成。
121+ T dataValue = AscendC::ReadGmBypassDCache<T>(addr + 1);
122+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(dstGm), 2 * dataValue);
123+ return;
124+ }
125+ }
126+}
127+```
128+ 
129+完整样例请参考[DataSyncBarrier样例](../../../../../../examples/01_simd_cpp_api/03_basic_api/05_sync_control/data_sync_barrier)。
Mdocs/zh/api/SIMD-API/basic_api/scalar_compute/scalar_compute.md+2-2
@@ -16,6 +16,6 @@
16 16 
17- **[Nop](Nop.md)**17- **[Nop](Nop.md)**
18 18 
19-- **[WriteGmByPassDCache\(ISASI\)](WriteGmByPassDCache_ISASI.md)**19+- **[WriteGmBypassDCache\(ISASI\)](WriteGmBypassDCache_ISASI.md)**
20 20 
21-- **[ReadGmByPassDCache\(ISASI\)](ReadGmByPassDCache_ISASI.md)**21+- **[ReadGmBypassDCache\(ISASI\)](ReadGmBypassDCache_ISASI.md)**
Mdocs/zh/api/SIMD-API/basic_api/sync_control/intra_core_sync/DataSyncBarrier_ISASI.md+5-5
@@ -59,22 +59,22 @@ __aicore__ inline void DataSyncBarrier()
59```cpp59```cpp
60 if (blockIdx == 0) {60 if (blockIdx == 0) {
61 // 先写入被依赖数据，再通过DataSyncBarrier等待GM访问完成。61 // 先写入被依赖数据，再通过DataSyncBarrier等待GM访问完成。
62- AscendC::WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm) + 1, DATA_VALUE);62+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm) + 1, DATA_VALUE);
63 // DataSyncBarrier<DDR>阻塞后续GM写，确保上一条GM写对其他核可见。63 // DataSyncBarrier<DDR>阻塞后续GM写，确保上一条GM写对其他核可见。
64 AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();64 AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
65 // 最后写入同步标记，block 1读到该标记后即可安全读取srcGm[1]。65 // 最后写入同步标记，block 1读到该标记后即可安全读取srcGm[1]。
66- AscendC::WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm), SYNC_FLAG);66+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(srcGm), SYNC_FLAG);
67 }67 }
68 68 
69 if (blockIdx == 1) {69 if (blockIdx == 1) {
70 while (true) {70 while (true) {
71 __gm__ T *addr = const_cast<__gm__ T *>(srcGlobal.GetPhyAddr());71 __gm__ T *addr = const_cast<__gm__ T *>(srcGlobal.GetPhyAddr());
72 // 轮询GM第0个元素，等待block 0写入同步标记。72 // 轮询GM第0个元素，等待block 0写入同步标记。
73- T flagValue = AscendC::ReadGmByPassDCache<T>(addr);73+ T flagValue = AscendC::ReadGmBypassDCache<T>(addr);
74 if (flagValue == SYNC_FLAG) {74 if (flagValue == SYNC_FLAG) {
75 // DataSyncBarrier保证同步标记之前的srcGm[1]写入已完成。75 // DataSyncBarrier保证同步标记之前的srcGm[1]写入已完成。
76- T dataValue = AscendC::ReadGmByPassDCache<T>(addr + 1);76+ T dataValue = AscendC::ReadGmBypassDCache<T>(addr + 1);
77- AscendC::WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T *>(dstGm), 2 * dataValue);77+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T *>(dstGm), 2 * dataValue);
78 return;78 return;
79 }79 }
80 }80 }
Mdocs/zh/api/appendix/api_pipeline_type_summary.md+2-3
@@ -227,10 +227,9 @@
227 227 
228| API接口名称 | 流水类型 |228| API接口名称 | 流水类型 |
229| --- | --- |229| --- | --- |
230-| ReadGmByPassDCache(ISASI) | PIPE_S |230+| ReadGmBypassDCache(ISASI) | PIPE_S |
231-| WriteGmByPassDCache(ISASI) | PIPE_S |231+| WriteGmBypassDCache(ISASI) | PIPE_S |
232| DataCachePreload | PIPE_S |232| DataCachePreload | PIPE_S |
233| DataCacheCleanAndInvalid | PIPE_S |233| DataCacheCleanAndInvalid | PIPE_S |
234| ICachePreLoad(ISASI) | PIPE_S |234| ICachePreLoad(ISASI) | PIPE_S |
235| GetICachePreloadStatus(ISASI) | PIPE_S |235| GetICachePreloadStatus(ISASI) | PIPE_S |
236- 
Mdocs/zh/api/appendix/deprecated_interface.md+10-0
@@ -1,5 +1,15 @@
1# 废弃接口<a name="ZH-CN_TOPIC_0000001724165365"></a>1# 废弃接口<a name="ZH-CN_TOPIC_0000001724165365"></a>
2 2 
3+## 基础API
4+ 
5+- [WriteGmByPassDCache\(ISASI\)](../SIMD-API/basic_api/deprecated_interface/WriteGmByPassDCache_ISASI_deprecated.md)接口
6+ 
7+ 此接口后续版本会废弃，请使用[WriteGmBypassDCache\(ISASI\)](../SIMD-API/basic_api/scalar_compute/WriteGmBypassDCache_ISASI.md)接口。
8+ 
9+- [ReadGmByPassDCache\(ISASI\)](../SIMD-API/basic_api/deprecated_interface/ReadGmByPassDCache_ISASI_deprecated.md)接口
10+ 
11+ 此接口后续版本会废弃，请使用[ReadGmBypassDCache\(ISASI\)](../SIMD-API/basic_api/scalar_compute/ReadGmBypassDCache_ISASI.md)接口。
12+ 
3## C API13## C API
4 14 
5### 数据搬运15### 数据搬运
Mdocs/zh/guide/programming_guide/advanced_programming/memory_model/cache_coherence.md+3-3
@@ -213,9 +213,9 @@ T6 Scalar.LOAD A -> 1
213 213 
214- 生产者通过Scalar写GM后，如果消费者需要从GM或其他路径读取该数据，生产者需要执行DCCI（例如`asc_dcci_single`）把Dirty数据写回GM。214- 生产者通过Scalar写GM后，如果消费者需要从GM或其他路径读取该数据，生产者需要执行DCCI（例如`asc_dcci_single`）把Dirty数据写回GM。
215- 消费者如果担心本核DCache中有旧副本，读取前需要执行DCCI使旧副本失效。215- 消费者如果担心本核DCache中有旧副本，读取前需要执行DCCI使旧副本失效。
216-- 如果算法本身不需要DCache缓存，可以使用[ReadGmByPassDCache](<../../../../api/SIMD-API/basic_api/scalar_compute/ReadGmByPassDCache_ISASI.md>)和[WriteGmByPassDCache](<../../../../api/SIMD-API/basic_api/scalar_compute/WriteGmByPassDCache_ISASI.md>)直接读写GM。216+- 如果算法本身不需要DCache缓存，可以使用[ReadGmBypassDCache](<../../../../api/SIMD-API/basic_api/scalar_compute/ReadGmBypassDCache_ISASI.md>)和[WriteGmBypassDCache](<../../../../api/SIMD-API/basic_api/scalar_compute/WriteGmBypassDCache_ISASI.md>)直接读写GM。
217 217 
218-还有一类问题来自Cache Line粒度。普通Scalar写后续以64B Cache Line为单位写回；多个核即使写的是同一条Cache Line中的不同字段，也可能在各自写回整条Cache Line时覆盖对方结果。此类场景可以通过让不同核操作的地址至少相隔Cache Line大小，或使用ReadGmByPassDCache和WriteGmByPassDCache降低风险。218+还有一类问题来自Cache Line粒度。普通Scalar写后续以64B Cache Line为单位写回；多个核即使写的是同一条Cache Line中的不同字段，也可能在各自写回整条Cache Line时覆盖对方结果。此类场景可以通过让不同核操作的地址至少相隔Cache Line大小，或使用ReadGmBypassDCache和WriteGmBypassDCache降低风险。
219 219 
220<!-- npu="950" id3 -->220<!-- npu="950" id3 -->
221#### Atomic访问导致的DCache缓存一致性221#### Atomic访问导致的DCache缓存一致性
@@ -329,7 +329,7 @@ T4 Scalar.LOAD A -> 1
329 329 
330Scalar写GM可以分为两类：330Scalar写GM可以分为两类：
331- cache: 普通Scalar写通过DCache修改GM，例如直接写`__gm__`地址；331- cache: 普通Scalar写通过DCache修改GM，例如直接写`__gm__`地址；
332-- no cache: 不经过DCache，直接向GM发起写操作，例如`WriteGmByPassDCache`等。332+- no cache: 不经过DCache，直接向GM发起写操作，例如`WriteGmBypassDCache`等。
333 333 
334按照访问顺序，常见需要软件维护一致性的场景可以分为RAW、WAW和WAR三类：334按照访问顺序，常见需要软件维护一致性的场景可以分为RAW、WAW和WAR三类：
335 335 
Mdocs/zh/guide/programming_guide/advanced_programming/memory_model/memory_consistency.md+2-2
@@ -407,7 +407,7 @@ T3 Scalar.STORE_GM_DCache(flag) = 1
407 407 
408##### Case 3：前后访问都是不经过DCache的访问408##### Case 3：前后访问都是不经过DCache的访问
409 409 
410-不经过DCache路径直接访问GM，不经过DCache，例如[asc_store_dev](../../../../api/SIMD-API/c_api/scalar_compute/scalar_store/asc_store_dev.md)、[WriteGmByPassDCache](../../../../api/SIMD-API/basic_api/scalar_compute/WriteGmByPassDCache_ISASI.md)、Scalar原子操作等接口。前后访问都不经过DCache时，无论访问同一地址还是不同地址，如果后序操作依赖前序访问已经完成，都需要用户插入DSB等待前序GM访问完成。410+不经过DCache路径直接访问GM，不经过DCache，例如[asc_store_dev](../../../../api/SIMD-API/c_api/scalar_compute/scalar_store/asc_store_dev.md)、[WriteGmBypassDCache](../../../../api/SIMD-API/basic_api/scalar_compute/WriteGmBypassDCache_ISASI.md)、Scalar原子操作等接口。前后访问都不经过DCache时，无论访问同一地址还是不同地址，如果后序操作依赖前序访问已经完成，都需要用户插入DSB等待前序GM访问完成。
411 411 
412同地址WAW场景中，如果期望最终结果以后一次写入为准，不能只依赖代码顺序：412同地址WAW场景中，如果期望最终结果以后一次写入为准，不能只依赖代码顺序：
413 413 
@@ -554,7 +554,7 @@ T2 STORE ready = 1
554 554 
555如果消费者访问GM时使用SIMT DCache，还需要同时结合SIMT DCache与GM的[缓存一致性](cache_coherence.md)要求处理旧副本。MEMBAR负责约束访问顺序，不负责自动失效消费者侧Cache副本。555如果消费者访问GM时使用SIMT DCache，还需要同时结合SIMT DCache与GM的[缓存一致性](cache_coherence.md)要求处理旧副本。MEMBAR负责约束访问顺序，不负责自动失效消费者侧Cache副本。
556 556 
557- 557+ 
558#### 线程间访问558#### 线程间访问
559 559 
560SIMT的不同线程之间，不论访问是否为同一地址，硬件不自动保证线程间内存一致性。一个线程写入的数据，如果要被另一个线程读取，需要用户自行插入同步。560SIMT的不同线程之间，不论访问是否为同一地址，硬件不自动保证线程间内存一致性。一个线程写入的数据，如果要被另一个线程读取，需要用户自行插入同步。
Mdocs/zh/guide/programming_guide/programming_model/ai_core_simd_programming/cpp_tensor_programming/cpp_tensor_programming_overview.md+2-2
@@ -347,7 +347,7 @@ __global__ __cube__ void matmul_kernel(__gm__ float* x, __gm__ float* y, __gm__
347 AscendC::LocalMemAllocator<AscendC::Hardware::L0A> l0aAllocator;347 AscendC::LocalMemAllocator<AscendC::Hardware::L0A> l0aAllocator;
348 AscendC::LocalMemAllocator<AscendC::Hardware::L0B> l0bAllocator;348 AscendC::LocalMemAllocator<AscendC::Hardware::L0B> l0bAllocator;
349 AscendC::LocalMemAllocator<AscendC::Hardware::L0C> l0cAllocator;349 AscendC::LocalMemAllocator<AscendC::Hardware::L0C> l0cAllocator;
350- 350+ 
351 AscendC::LocalTensor<half> a_l1_buf = l1Allocator.Alloc<half, BLOCK_M * BLOCK_K>();351 AscendC::LocalTensor<half> a_l1_buf = l1Allocator.Alloc<half, BLOCK_M * BLOCK_K>();
352 AscendC::LocalTensor<half> b_l1_buf = l1Allocator.Alloc<half, BLOCK_K * BLOCK_N>();352 AscendC::LocalTensor<half> b_l1_buf = l1Allocator.Alloc<half, BLOCK_K * BLOCK_N>();
353 AscendC::LocalTensor<half> l0a_buf = l0aAllocator.Alloc<half, BLOCK_M * BLOCK_K>();353 AscendC::LocalTensor<half> l0a_buf = l0aAllocator.Alloc<half, BLOCK_M * BLOCK_K>();
@@ -943,7 +943,7 @@ __global__ __mix__ void mix_kernel(__gm__ float* x, __gm__ float* y, __gm__ floa
943</thead>943</thead>
944<tbody><tr id="row1298763513418"><td class="cellrowborder" valign="top" width="29.630000000000003%" headers="mcps1.2.4.1.1 "><p id="p15987173524114"><a name="p15987173524114"></a><a name="p15987173524114"></a>基础API &gt; 标量计算</p>944<tbody><tr id="row1298763513418"><td class="cellrowborder" valign="top" width="29.630000000000003%" headers="mcps1.2.4.1.1 "><p id="p15987173524114"><a name="p15987173524114"></a><a name="p15987173524114"></a>基础API &gt; 标量计算</p>
945</td>945</td>
946-<td class="cellrowborder" valign="top" width="44.519999999999996%" headers="mcps1.2.4.1.2 "><p id="p12167195363614"><a name="p12167195363614"></a><a name="p12167195363614"></a>GetBitCount、CountLeadingZero、CountBitsCntSameAsSignBit、GetSFFValue、Cast（float转half、int32_t）、Cast（float转bfloat16_t）、Cast（多类型转float）、Nop、GetUintDivMagicAndShift、WriteGmByPassDCache、ReadGmByPassDCache</p>946+<td class="cellrowborder" valign="top" width="44.519999999999996%" headers="mcps1.2.4.1.2 "><p id="p12167195363614"><a name="p12167195363614"></a><a name="p12167195363614"></a>GetBitCount、CountLeadingZero、CountBitsCntSameAsSignBit、GetSFFValue、Cast（float转half、int32_t）、Cast（float转bfloat16_t）、Cast（多类型转float）、Nop、GetUintDivMagicAndShift、WriteGmBypassDCache、ReadGmBypassDCache</p>
947</td>947</td>
948<td class="cellrowborder" valign="top" width="25.85%" headers="mcps1.2.4.1.3 "><p id="p298711352411"><a name="p298711352411"></a><a name="p298711352411"></a>-</p>948<td class="cellrowborder" valign="top" width="25.85%" headers="mcps1.2.4.1.3 "><p id="p298711352411"><a name="p298711352411"></a><a name="p298711352411"></a>-</p>
949</td>949</td>
Mdocs/zh/guide/programming_guide/programming_model/ai_core_simd_programming/cpp_tensor_programming/static_tensor_programming.md+1-1
@@ -727,7 +727,7 @@
727</thead>727</thead>
728<tbody><tr id="row1298763513418"><td class="cellrowborder" valign="top" width="29.630000000000003%" headers="mcps1.2.4.1.1 "><p id="p15987173524114"><a name="p15987173524114"></a><a name="p15987173524114"></a>基础API &gt; 标量计算</p>728<tbody><tr id="row1298763513418"><td class="cellrowborder" valign="top" width="29.630000000000003%" headers="mcps1.2.4.1.1 "><p id="p15987173524114"><a name="p15987173524114"></a><a name="p15987173524114"></a>基础API &gt; 标量计算</p>
729</td>729</td>
730-<td class="cellrowborder" valign="top" width="44.519999999999996%" headers="mcps1.2.4.1.2 "><p id="p12167195363614"><a name="p12167195363614"></a><a name="p12167195363614"></a>GetBitCount、CountLeadingZero、CountBitsCntSameAsSignBit、GetSFFValue、Cast（float转half、int32_t）、Cast（float转bfloat16_t）、Cast（多类型转float）、Nop、GetUintDivMagicAndShift、WriteGmByPassDCache、ReadGmByPassDCache</p>730+<td class="cellrowborder" valign="top" width="44.519999999999996%" headers="mcps1.2.4.1.2 "><p id="p12167195363614"><a name="p12167195363614"></a><a name="p12167195363614"></a>GetBitCount、CountLeadingZero、CountBitsCntSameAsSignBit、GetSFFValue、Cast（float转half、int32_t）、Cast（float转bfloat16_t）、Cast（多类型转float）、Nop、GetUintDivMagicAndShift、WriteGmBypassDCache、ReadGmBypassDCache</p>
731</td>731</td>
732<td class="cellrowborder" valign="top" width="25.85%" headers="mcps1.2.4.1.3 "><p id="p298711352411"><a name="p298711352411"></a><a name="p298711352411"></a>-</p>732<td class="cellrowborder" valign="top" width="25.85%" headers="mcps1.2.4.1.3 "><p id="p298711352411"><a name="p298711352411"></a><a name="p298711352411"></a>-</p>
733</td>733</td>
Mdocs/zh/vitepress/docs/.vitepress/theme/outline-index.json+1-1
此文件变更行数或变更字符数较多，你可以直接 查看源码
Mexamples/01_simd_cpp_api/03_basic_api/05_sync_control/data_sync_barrier/README.md+1-1
@@ -40,7 +40,7 @@
40下面按步骤介绍本样例中 `DataSyncBarrier` 的使用场景：40下面按步骤介绍本样例中 `DataSyncBarrier` 的使用场景：
41 41 
421. 系统中有两个 AIV 核，分别记为核 0 和核 1。GM 中的两个变量 `x` 和 `y` 初始值都为 1。421. 系统中有两个 AIV 核，分别记为核 0 和核 1。GM 中的两个变量 `x` 和 `y` 初始值都为 1。
43-2. 核 0 先通过标量流水接口`WriteGmByPassDCache`向 `srcGm[1]` 写入 `x=7`。43+2. 核 0 先通过标量流水接口`WriteGmBypassDCache`向 `srcGm[1]` 写入 `x=7`。
443. 核 0 再插入 `DataSyncBarrier<AscendC::MemDsbT::DDR>()`，等待前一次 GM 写操作完成。443. 核 0 再插入 `DataSyncBarrier<AscendC::MemDsbT::DDR>()`，等待前一次 GM 写操作完成。
454. 核 0 随后向 `srcGm[0]` 写入 `y=6`。454. 核 0 随后向 `srcGm[0]` 写入 `y=6`。
465. 核 1 持续轮询 `srcGm[0]`，直到读到 `y=6`，再读取 `srcGm[1]`，并将 `2 * x` 写入输出。465. 核 1 持续轮询 `srcGm[0]`，直到读到 `y=6`，再读取 `srcGm[1]`，并将 `2 * x` 写入输出。
Mexamples/01_simd_cpp_api/03_basic_api/05_sync_control/data_sync_barrier/README_en.md+1-1
@@ -40,7 +40,7 @@ This example introduces the invocation of DataSyncBarrier. This interface blocks
40The following describes the usage scenario of `DataSyncBarrier` in this example step by step:40The following describes the usage scenario of `DataSyncBarrier` in this example step by step:
41 41 
421. The system has two AIV cores, denoted as Core 0 and Core 1. Two variables `x` and `y` in GM both have an initial value of 1.421. The system has two AIV cores, denoted as Core 0 and Core 1. Two variables `x` and `y` in GM both have an initial value of 1.
43-2. Core 0 first writes `x=7` to `srcGm[1]` through the scalar pipeline interface `WriteGmByPassDCache`.43+2. Core 0 first writes `x=7` to `srcGm[1]` through the scalar pipeline interface `WriteGmBypassDCache`.
443. Core 0 then inserts `DataSyncBarrier<AscendC::MemDsbT::DDR>()` to wait for the previous GM write operation to complete.443. Core 0 then inserts `DataSyncBarrier<AscendC::MemDsbT::DDR>()` to wait for the previous GM write operation to complete.
454. Core 0 subsequently writes `y=6` to `srcGm[0]`.454. Core 0 subsequently writes `y=6` to `srcGm[0]`.
465. Core 1 continuously polls `srcGm[0]` until it reads `y=6`, then reads `srcGm[1]` and writes `2 * x` to the output.465. Core 1 continuously polls `srcGm[0]` until it reads `y=6`, then reads `srcGm[1]` and writes `2 * x` to the output.
Mexamples/01_simd_cpp_api/03_basic_api/05_sync_control/data_sync_barrier/data_sync_barrier.asc+5-5
@@ -31,22 +31,22 @@ __global__ __vector__ void kernel_data_sync_barrier(__gm__ uint8_t* srcGm, __gm_
31 31 
32 if (blockIdx == 0) {32 if (blockIdx == 0) {
33 // 先写入被依赖数据，再通过DataSyncBarrier等待DDR访问完成。33 // 先写入被依赖数据，再通过DataSyncBarrier等待DDR访问完成。
34- AscendC::WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(srcGm) + 1, DATA_VALUE);34+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T*>(srcGm) + 1, DATA_VALUE);
35 // DataSyncBarrier<DDR>阻塞后续GM写，确保上一条GM写对其他核可见。35 // DataSyncBarrier<DDR>阻塞后续GM写，确保上一条GM写对其他核可见。
36 AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();36 AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>();
37 // 最后写入同步标记，block 1读到该标记后即可安全读取srcGm[1]。37 // 最后写入同步标记，block 1读到该标记后即可安全读取srcGm[1]。
38- AscendC::WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(srcGm), SYNC_FLAG);38+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T*>(srcGm), SYNC_FLAG);
39 }39 }
40 40 
41 if (blockIdx == 1) {41 if (blockIdx == 1) {
42 while (true) {42 while (true) {
43 __gm__ T* addr = const_cast<__gm__ T*>(srcGlobal.GetPhyAddr());43 __gm__ T* addr = const_cast<__gm__ T*>(srcGlobal.GetPhyAddr());
44 // 轮询GM第0个元素，等待block 0写入同步标记。44 // 轮询GM第0个元素，等待block 0写入同步标记。
45- T flagValue = AscendC::ReadGmByPassDCache<T>(addr);45+ T flagValue = AscendC::ReadGmBypassDCache<T>(addr);
46 if (flagValue == SYNC_FLAG) {46 if (flagValue == SYNC_FLAG) {
47 // DataSyncBarrier保证同步标记之前的srcGm[1]写入已完成。47 // DataSyncBarrier保证同步标记之前的srcGm[1]写入已完成。
48- T dataValue = AscendC::ReadGmByPassDCache<T>(addr + 1);48+ T dataValue = AscendC::ReadGmBypassDCache<T>(addr + 1);
49- AscendC::WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(dstGm), 2 * dataValue);49+ AscendC::WriteGmBypassDCache<T>(reinterpret_cast<__gm__ T*>(dstGm), 2 * dataValue);
50 return;50 return;
51 }51 }
52 }52 }
Mexamples/01_simd_cpp_api/03_basic_api/06_atomic/set_atomic_deterministic_computation/README.md+1-1
@@ -217,7 +217,7 @@ Scalar单元访问GM上的信号量，存在两种访问方式：
217 使用 GlobalTensor的成员函数GetValue 和 SetValue 进行操作。这种情况下，开发者需要手动调用[DataCacheCleanAndInvalid](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_0177.html)接口，以确保多核间数据的一致性。217 使用 GlobalTensor的成员函数GetValue 和 SetValue 进行操作。这种情况下，开发者需要手动调用[DataCacheCleanAndInvalid](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_0177.html)接口，以确保多核间数据的一致性。
218 218 
2192. **不通过DCache 访问**2192. **不通过DCache 访问**
220- 使用[WriteGmByPassDCache](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_00089.html) 和 ReadGmByPassDCache。这种方式无需额外操作即可保证多核间数据的一致性。220+ 使用[WriteGmBypassDCache](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_00089.html) 和 ReadGmBypassDCache。这种方式无需额外操作即可保证多核间数据的一致性。
221 221 
222两种方案在性能上的差异：不经过DCache，性能会较差，但是如果读写GM数据较少，可以考虑使用不经过DCache的方法。222两种方案在性能上的差异：不经过DCache，性能会较差，但是如果读写GM数据较少，可以考虑使用不经过DCache的方法。
223核间同步方案中也需要与核内同步配合使用，现将三处核内同步作用说明如下：223核间同步方案中也需要与核内同步配合使用，现将三处核内同步作用说明如下：
Mexamples/01_simd_cpp_api/03_basic_api/06_atomic/set_atomic_deterministic_computation/README_en.md+1-1
@@ -217,7 +217,7 @@ The Scalar unit accesses semaphores on GM through two access methods:
217 Use the GetValue and SetValue member functions of GlobalTensor. In this case, developers must manually call the [DataCacheCleanAndInvalid](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_0177.html) API to ensure data consistency across multiple cores.217 Use the GetValue and SetValue member functions of GlobalTensor. In this case, developers must manually call the [DataCacheCleanAndInvalid](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_0177.html) API to ensure data consistency across multiple cores.
218 218 
2192. **Access bypassing DCache**2192. **Access bypassing DCache**
220- Use [WriteGmByPassDCache](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_00089.html) and ReadGmByPassDCache. This method ensures multi-core data consistency without additional operations.220+ Use [WriteGmBypassDCache](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/API/ascendcopapi/atlasascendc_api_07_00089.html) and ReadGmBypassDCache. This method ensures multi-core data consistency without additional operations.
221 221 
222Performance difference between the two schemes: bypassing DCache has relatively poorer performance, but if the amount of data read from or written to GM is small, the bypass-DCache method can be considered.222Performance difference between the two schemes: bypassing DCache has relatively poorer performance, but if the amount of data read from or written to GM is small, the bypass-DCache method can be considered.
223The inter-core synchronization scheme also requires intra-core synchronization. The roles of the three intra-core synchronization points are described as follows:223The inter-core synchronization scheme also requires intra-core synchronization. The roles of the three intra-core synchronization points are described as follows:
Mexamples/01_simd_cpp_api/03_basic_api/06_atomic/set_atomic_deterministic_computation/set_atomic_add_multi_aic.h+9-9
@@ -60,15 +60,15 @@ public:
60 // 在通知下一个核之前，等待当前核的任务完成60 // 在通知下一个核之前，等待当前核的任务完成
61 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);61 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);
62 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);62 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);
63- AscendC::WriteGmByPassDCache<int32_t>(addr + blockIdx * 32, 1);63+ AscendC::WriteGmBypassDCache<int32_t>(addr + blockIdx * 32, 1);
64 } else if (blockIdx == 1) {64 } else if (blockIdx == 1) {
65 int32_t preblockIdx = 3;65 int32_t preblockIdx = 3;
66 66 
67 // 核 1 等待核 3，直到核 3 数据已完成累加67 // 核 1 等待核 3，直到核 3 数据已完成累加
68 while (true) {68 while (true) {
69- int32_t value = AscendC::ReadGmByPassDCache<int32_t>(addr + preblockIdx * 32);69+ int32_t value = AscendC::ReadGmBypassDCache<int32_t>(addr + preblockIdx * 32);
70 if (value == 1) {70 if (value == 1) {
71- AscendC::WriteGmByPassDCache<int32_t>(addr + preblockIdx * 32, 0);71+ AscendC::WriteGmBypassDCache<int32_t>(addr + preblockIdx * 32, 0);
72 break;72 break;
73 }73 }
74 }74 }
@@ -89,9 +89,9 @@ public:
89 89 
90 // 核 2 等待核 0，直到核 0 数据已完成累加90 // 核 2 等待核 0，直到核 0 数据已完成累加
91 while (true) {91 while (true) {
92- int32_t value = AscendC::ReadGmByPassDCache<int32_t>(addr + preblockIdx * 32);92+ int32_t value = AscendC::ReadGmBypassDCache<int32_t>(addr + preblockIdx * 32);
93 if (value == 1) {93 if (value == 1) {
94- AscendC::WriteGmByPassDCache<int32_t>(addr + preblockIdx * 32, 0);94+ AscendC::WriteGmBypassDCache<int32_t>(addr + preblockIdx * 32, 0);
95 break;95 break;
96 }96 }
97 }97 }
@@ -111,15 +111,15 @@ public:
111 // 在通知下一个核之前，等待当前核的任务完成111 // 在通知下一个核之前，等待当前核的任务完成
112 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);112 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);
113 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);113 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);
114- AscendC::WriteGmByPassDCache<int32_t>(addr + blockIdx * 32, 1);114+ AscendC::WriteGmBypassDCache<int32_t>(addr + blockIdx * 32, 1);
115 } else if (blockIdx == 3) {115 } else if (blockIdx == 3) {
116 int32_t preblockIdx = 2;116 int32_t preblockIdx = 2;
117 117 
118 // 核 3 等待核 2，直到核 2 数据已完成累加118 // 核 3 等待核 2，直到核 2 数据已完成累加
119 while (true) {119 while (true) {
120- int32_t value = AscendC::ReadGmByPassDCache<int32_t>(addr + preblockIdx * 32);120+ int32_t value = AscendC::ReadGmBypassDCache<int32_t>(addr + preblockIdx * 32);
121 if (value == 1) {121 if (value == 1) {
122- AscendC::WriteGmByPassDCache<int32_t>(addr + preblockIdx * 32, 0);122+ AscendC::WriteGmBypassDCache<int32_t>(addr + preblockIdx * 32, 0);
123 break;123 break;
124 }124 }
125 }125 }
@@ -140,7 +140,7 @@ public:
140 // 在通知下一个核之前，等待当前核的任务完成140 // 在通知下一个核之前，等待当前核的任务完成
141 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);141 AscendC::SetFlag<AscendC::HardEvent::MTE3_S>(0);
142 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);142 AscendC::WaitFlag<AscendC::HardEvent::MTE3_S>(0);
143- AscendC::WriteGmByPassDCache<int32_t>(addr + blockIdx * 32, 1);143+ AscendC::WriteGmBypassDCache<int32_t>(addr + blockIdx * 32, 1);
144 }144 }
145 }145 }
146 146 
Mexamples/01_simd_cpp_api/03_basic_api/09_utils/README.md+1-1
@@ -11,4 +11,4 @@
11| [ctrl_spr](./ctrl_spr) | 本样例基于SetCtrlSpr、GetCtrlSpr、ResetCtrlSpr接口实现对CTRL寄存器（控制寄存器）的特定比特位的设置、读取和重置，并验证非饱和模式是否正确。 | Ascend 950PR/Ascend 950DT |11| [ctrl_spr](./ctrl_spr) | 本样例基于SetCtrlSpr、GetCtrlSpr、ResetCtrlSpr接口实现对CTRL寄存器（控制寄存器）的特定比特位的设置、读取和重置，并验证非饱和模式是否正确。 | Ascend 950PR/Ascend 950DT |
12| [get_ub_size](./get_ub_size) | 本样例展示GetUBSizeInBytes和GetRuntimeUBSize接口使用方法，用于获取用户最大可使用的UB大小。 | Ascend 950PR/Ascend 950DT |12| [get_ub_size](./get_ub_size) | 本样例展示GetUBSizeInBytes和GetRuntimeUBSize接口使用方法，用于获取用户最大可使用的UB大小。 | Ascend 950PR/Ascend 950DT |
13| [gm_by_dcache](./gm_by_dcache) | 本样例介绍经过DCache（Data Cache）从GM地址读取和写入数据。 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |13| [gm_by_dcache](./gm_by_dcache) | 本样例介绍经过DCache（Data Cache）从GM地址读取和写入数据。 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
14-| [gm_by_pass_dcache](./gm_by_pass_dcache) | 本样例基于ReadGmByPassDCache和WriteGmByPassDCache接口，实现不经过DCache从GM读取和写入数据。 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |14+| [gm_by_pass_dcache](./gm_by_pass_dcache) | 本样例基于ReadGmBypassDCache和WriteGmBypassDCache接口，实现不经过DCache从GM读取和写入数据。 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
Mexamples/01_simd_cpp_api/03_basic_api/09_utils/README_en.md+1-1
@@ -11,4 +11,4 @@ This directory contains API samples related to utilities. The samples are based
11| [ctrl_spr](./ctrl_spr) | This sample demonstrates setting, reading, and resetting specific bit fields of the CTRL register (control register) using SetCtrlSpr, GetCtrlSpr, and ResetCtrlSpr interfaces, and verifies whether non-saturation mode is working correctly. | Ascend 950PR/Ascend 950DT |11| [ctrl_spr](./ctrl_spr) | This sample demonstrates setting, reading, and resetting specific bit fields of the CTRL register (control register) using SetCtrlSpr, GetCtrlSpr, and ResetCtrlSpr interfaces, and verifies whether non-saturation mode is working correctly. | Ascend 950PR/Ascend 950DT |
12| [get_ub_size](./get_ub_size) | This sample demonstrates the usage of GetUBSizeInBytes and GetRuntimeUBSize interfaces to get the maximum user-available UB size. | Ascend 950PR/Ascend 950DT |12| [get_ub_size](./get_ub_size) | This sample demonstrates the usage of GetUBSizeInBytes and GetRuntimeUBSize interfaces to get the maximum user-available UB size. | Ascend 950PR/Ascend 950DT |
13| [gm_by_dcache](./gm_by_dcache) | This sample demonstrates reading data from and writing data to GM with going through DCache. | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |13| [gm_by_dcache](./gm_by_dcache) | This sample demonstrates reading data from and writing data to GM with going through DCache. | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |
14-| [gm_by_pass_dcache](./gm_by_pass_dcache) | This sample demonstrates reading data from and writing data to GM without going through DCache, based on the ReadGmByPassDCache and WriteGmByPassDCache interfaces. | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |14+| [gm_by_pass_dcache](./gm_by_pass_dcache) | This sample demonstrates reading data from and writing data to GM without going through DCache, based on the ReadGmBypassDCache and WriteGmBypassDCache interfaces. | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |
Mexamples/01_simd_cpp_api/03_basic_api/09_utils/gm_by_pass_dcache/README.md+4-4
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例组合调用 `ReadGmByPassDCache` 和 `WriteGmByPassDCache` 接口，实现不经过DCache从GM地址读取和写入数据。本样例会将读取的值加100后写入到输出地址。5+本样例组合调用 `ReadGmBypassDCache` 和 `WriteGmBypassDCache` 接口，实现不经过DCache从GM地址读取和写入数据。本样例会将读取的值加100后写入到输出地址。
6 6 
7当多核操作GM地址时，如果数据无法对齐到Cache Line，经过DCache的方式下，由于按照Cache Line大小进行读写，会导致多核数据随机覆盖的问题。此时，可以采用不经过DCache直接读写GM地址的方式，从而避免上述随机覆盖的问题。7当多核操作GM地址时，如果数据无法对齐到Cache Line，经过DCache的方式下，由于按照Cache Line大小进行读写，会导致多核数据随机覆盖的问题。此时，可以采用不经过DCache直接读写GM地址的方式，从而避免上述随机覆盖的问题。
8 8 
@@ -30,9 +30,9 @@
30## 样例描述30## 样例描述
31 31 
32- 样例功能： 32- 样例功能：
33- 组合调用 `ReadGmByPassDCache` 和 `WriteGmByPassDCache`，实现：33+ 组合调用 `ReadGmBypassDCache` 和 `WriteGmBypassDCache`，实现：
34- 1. 使用 `ReadGmByPassDCache` 从 GM 地址逐个读取int32_t类型元素34+ 1. 使用 `ReadGmBypassDCache` 从 GM 地址逐个读取int32_t类型元素
35- 2. 使用 `WriteGmByPassDCache` 将读出的值加100后，写入到输出GM地址35+ 2. 使用 `WriteGmBypassDCache` 将读出的值加100后，写入到输出GM地址
36 36 
37- 样例规格：37- 样例规格：
38 <table>38 <table>
Mexamples/01_simd_cpp_api/03_basic_api/09_utils/gm_by_pass_dcache/README_en.md+4-4
@@ -2,7 +2,7 @@
2 2 
3## Overview3## Overview
4 4 
5-This example combines calls to the `ReadGmByPassDCache` and `WriteGmByPassDCache` interfaces to read and write data from GM addresses without going through DCache. This example adds 100 to the read values and writes them to the output address.5+This example combines calls to the `ReadGmBypassDCache` and `WriteGmBypassDCache` interfaces to read and write data from GM addresses without going through DCache. This example adds 100 to the read values and writes them to the output address.
6 6 
7When multiple cores operate on GM addresses, if the data cannot be aligned to Cache Line boundaries, the DCache path reads and writes in Cache Line sizes, which causes random data overwriting between cores. In this case, the approach of reading and writing GM addresses directly without going through DCache can be adopted to avoid the random overwriting issue.7When multiple cores operate on GM addresses, if the data cannot be aligned to Cache Line boundaries, the DCache path reads and writes in Cache Line sizes, which causes random data overwriting between cores. In this case, the approach of reading and writing GM addresses directly without going through DCache can be adopted to avoid the random overwriting issue.
8 8 
@@ -30,9 +30,9 @@ When multiple cores operate on GM addresses, if the data cannot be aligned to Ca
30## Example Description30## Example Description
31 31 
32- Example function: 32- Example function:
33- Combine calls to `ReadGmByPassDCache` and `WriteGmByPassDCache` to implement:33+ Combine calls to `ReadGmBypassDCache` and `WriteGmBypassDCache` to implement:
34- 1. Use `ReadGmByPassDCache` to read int32_t elements one by one from GM addresses34+ 1. Use `ReadGmBypassDCache` to read int32_t elements one by one from GM addresses
35- 2. Use `WriteGmByPassDCache` to write the read values plus 100 to the output GM addresses35+ 2. Use `WriteGmBypassDCache` to write the read values plus 100 to the output GM addresses
36 36 
37- Example specifications:37- Example specifications:
38 <table>38 <table>
Mexamples/01_simd_cpp_api/03_basic_api/09_utils/gm_by_pass_dcache/gm_by_pass_dcache.asc+3-3
@@ -10,7 +10,7 @@
10 10 
11/* !11/* !
12 * \file gm_by_pass_dcache.asc12 * \file gm_by_pass_dcache.asc
13- * \brief 基于WriteGmByPassDCache和ReadGmByPassDCache接口，实现不经过DCache从GM地址读取和写入数据13+ * \brief 基于WriteGmBypassDCache和ReadGmBypassDCache接口，实现不经过DCache从GM地址读取和写入数据
14 */14 */
15 15 
16#include "acl/acl.h"16#include "acl/acl.h"
@@ -33,11 +33,11 @@ public:
33 __aicore__ inline void Process()33 __aicore__ inline void Process()
34 {34 {
35 __gm__ T* srcAddr = const_cast<__gm__ T*>(srcGlobal.GetPhyAddr());35 __gm__ T* srcAddr = const_cast<__gm__ T*>(srcGlobal.GetPhyAddr());
36- T value = AscendC::ReadGmByPassDCache<T>(srcAddr);36+ T value = AscendC::ReadGmBypassDCache<T>(srcAddr);
37 // 同步指令：用于阻塞后续的指令执行，直到所有之前的内存访问指令（需要等待的内存位置可通过参数控制）执行结束37 // 同步指令：用于阻塞后续的指令执行，直到所有之前的内存访问指令（需要等待的内存位置可通过参数控制）执行结束
38 AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>(); // DDR，等待GM访问指令38 AscendC::DataSyncBarrier<AscendC::MemDsbT::DDR>(); // DDR，等待GM访问指令
39 __gm__ T* dstAddr = const_cast<__gm__ T*>(dstGlobal.GetPhyAddr());39 __gm__ T* dstAddr = const_cast<__gm__ T*>(dstGlobal.GetPhyAddr());
40- AscendC::WriteGmByPassDCache<T>(dstAddr, value + ADD_VALUE);40+ AscendC::WriteGmBypassDCache<T>(dstAddr, value + ADD_VALUE);
41 }41 }
42 42 
43private:43private:
Mimpl/basic_api/kernel_operator_scalar_intf_impl.h+16-2
@@ -81,15 +81,29 @@ __aicore__ inline U ScalarCast(T valueIn)
81 81 
82#if __NPU_ARCH__ == 2201 || (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102)82#if __NPU_ARCH__ == 2201 || (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102)
83template <typename T>83template <typename T>
84+[[deprecated("Deprecated since 9.2.0, will be removed after 2027/12/30, Use WriteGmBypassDCache instead.")]]
84__aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value)85__aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value)
85{86{
86- return WriteGmByPassDCacheImpl(addr, value);87+ return WriteGmBypassDCacheImpl(addr, value);
87}88}
88 89 
89template <typename T>90template <typename T>
91+[[deprecated("Deprecated since 9.2.0, will be removed after 2027/12/30, Use ReadGmBypassDCache instead.")]]
90__aicore__ inline T ReadGmByPassDCache(__gm__ T* addr)92__aicore__ inline T ReadGmByPassDCache(__gm__ T* addr)
91{93{
92- return ReadGmByPassDCacheImpl(addr);94+ return ReadGmBypassDCacheImpl(addr);
95+}
96+ 
97+template <typename T>
98+__aicore__ inline void WriteGmBypassDCache(__gm__ T* addr, T value)
99+{
100+ return WriteGmBypassDCacheImpl(addr, value);
101+}
atomgit-bot
atomgit-botatomgit-bot22 天前

🔴 Critical

变更行 97 新增了 __aicore__ inline void WriteGmBypassDCache(__gm__ T* addr, T value)，参数列表使用了模板参数 T，但函数前缺少 template <typename T> 声明（对比紧邻的第 102-103 行 ReadGmBypassDCache 有完整的 template <typename T>）。在 namespace AscendC 作用域内不存在名为 T 的类型，编译器解析到该函数声明时即报 'T' does not name a type / use of undeclared identifier 'T'。该 impl 文件被公共头 include/basic_api/kernel_operator_scalar_intf.h（第 104-106 行，#if defined(__NPU_ARCH__) 时 include）引入，且所在代码块在 NPU_ARCH == 2201/3510/5102 下必然参与编译，因此任何在这三个架构上编译的算子代码都会直接编译失败（构建中断）。修复：在该函数前补上 template <typename T>，与公共头第 97-98 行的模板声明保持一致。

建议：在 WriteGmBypassDCache 函数定义前补加 template <typename T>，与公共头文件中的声明（第 97-98 行）及相邻 ReadGmBypassDCache 的写法一致。

likedislike
不准确？
102+ 
103+template <typename T>
104+__aicore__ inline T ReadGmBypassDCache(__gm__ T* addr)
105+{
106+ return ReadGmBypassDCacheImpl(addr);
93}107}
94#endif108#endif
95} // namespace AscendC109} // namespace AscendC
Mimpl/basic_api/kernel_scalar.h+15-4
@@ -116,13 +116,24 @@ __aicore__ inline int64_t GetSFFValueImpl(uint64_t valueIn)
116 116 
117#if __NPU_ARCH__ == 2201 || (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102)117#if __NPU_ARCH__ == 2201 || (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102)
118template <typename T>118template <typename T>
119-__aicore__ inline void WriteGmByPassDCacheImpl(__gm__ T* addr, T value)119+__aicore__ inline void WriteGmBypassDCacheImpl(__gm__ T* addr, T value)
120{120{
121 static_assert(121 static_assert(
122 SupportType<T, int8_t, uint8_t, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t>(),122 SupportType<T, int8_t, uint8_t, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t>(),
123- "WriteGmByPassDCache only support int8_t/uint8_t/int16_t/uint16_t/int32_t/uint32_t/int64_t/uint64_t "123+ "WriteGmBypassDCache only support int8_t/uint8_t/int16_t/uint16_t/int32_t/uint32_t/int64_t/uint64_t "
124 "data type on current device!");124 "data type on current device!");
125 125 
126+#if __NPU_ARCH__ == 2201
127+ // The write only takes effect in the first 32 bytes of each 128-byte aligned range.
128+ ASCENDC_DEBUG_ASSERT(
129+ (static_cast<uint32_t>(reinterpret_cast<uint64_t>(addr) % 128) < 32),
130+ KERNEL_LOG_INTERNAL(
131+ KERNEL_ERROR,
132+ "For WriteGmBypassDCache, current addr offset in a 128-byte aligned range is %u, which should be in "
133+ "range [0, 32).",
134+ static_cast<uint32_t>(reinterpret_cast<uint64_t>(addr) % 128)));
135+#endif
136+ 
126 if constexpr (SupportBytes<T, 8>()) {137 if constexpr (SupportBytes<T, 8>()) {
127 st_dev(*(reinterpret_cast<uint64_t*>(&value)), reinterpret_cast<__gm__ uint64_t*>(addr), 0);138 st_dev(*(reinterpret_cast<uint64_t*>(&value)), reinterpret_cast<__gm__ uint64_t*>(addr), 0);
128 } else if constexpr (SupportBytes<T, 4>()) {139 } else if constexpr (SupportBytes<T, 4>()) {
@@ -135,11 +146,11 @@ __aicore__ inline void WriteGmByPassDCacheImpl(__gm__ T* addr, T value)
135}146}
136 147 
137template <typename T>148template <typename T>
138-__aicore__ inline T ReadGmByPassDCacheImpl(__gm__ T* addr)149+__aicore__ inline T ReadGmBypassDCacheImpl(__gm__ T* addr)
139{150{
140 static_assert(151 static_assert(
141 SupportType<T, int8_t, uint8_t, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t>(),152 SupportType<T, int8_t, uint8_t, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t>(),
GRJ_XIDUO
GRJ_XIDUOGRJ_XIDUO19 天前

153行，告警信息中仍然是旧接口名

likedislike
142- "ReadGmByPassDCache only support int8_t/uint8_t/int16_t/uint16_t/int32_t/uint32_t/int64_t/uint64_t "153+ "ReadGmBypassDCache only support int8_t/uint8_t/int16_t/uint16_t/int32_t/uint32_t/int64_t/uint64_t "
143 "data type on current device!");154 "data type on current device!");
144 155 
145 if constexpr (SupportBytes<T, 8>()) {156 if constexpr (SupportBytes<T, 8>()) {
Minclude/basic_api/kernel_operator_scalar_intf.h+8-0
@@ -90,10 +90,18 @@ __aicore__ constexpr inline float Cast(const T& bVal);
90#endif90#endif
91 91 
92template <typename T>92template <typename T>
93+[[deprecated("Deprecated since 9.2.0, will be removed after 2027/12/30, Use WriteGmBypassDCache instead.")]]
93__aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value);94__aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value);
94 95 
95template <typename T>96template <typename T>
97+[[deprecated("Deprecated since 9.2.0, will be removed after 2027/12/30, Use ReadGmBypassDCache instead.")]]
96__aicore__ inline T ReadGmByPassDCache(__gm__ T* addr);98__aicore__ inline T ReadGmByPassDCache(__gm__ T* addr);
99+ 
100+template <typename T>
101+__aicore__ inline void WriteGmBypassDCache(__gm__ T* addr, T value);
102+ 
103+template <typename T>
104+__aicore__ inline T ReadGmBypassDCache(__gm__ T* addr);
97} // namespace AscendC105} // namespace AscendC
98 106 
99#if defined(__NPU_ARCH__)107#if defined(__NPU_ARCH__)
Mtests/api/basic_api/ascendc_case_ascend950pr_9599/ascendc_case_ascend950pr_9599_aiv_basic/test_global_memory_access.cpp+53-42
@@ -17,47 +17,50 @@ using AscendC::Reg::MaskReg;
17using AscendC::Reg::RegTensor;17using AscendC::Reg::RegTensor;
18using AscendC::Reg::UpdateMask;18using AscendC::Reg::UpdateMask;
19 19 
20-template <typename T, uint8_t mode>20+#define DEFINE_GM_ACCESS_CASE(READ_API, WRITE_API, TAG) \
21-class KernelGmAccess {21+ template <typename T, uint8_t Mode> \
22-public:22+ class KernelGmAccess##TAG { \
23- __aicore__ inline KernelGmAccess() {}23+ public: \
24- 24+ __aicore__ inline KernelGmAccess##TAG() {} \
25- __aicore__ inline void Process(GM_ADDR dst0Gm, GM_ADDR src0Gm, uint32_t calcount)25+ __aicore__ inline void Process(GM_ADDR dst0Gm, GM_ADDR src0Gm, uint32_t calcount) \
26- {26+ { \
27- if (mode == 0) {27+ __gm__ T* dstAddr = reinterpret_cast<__gm__ T*>(dst0Gm); \
28- T value = ReadGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(src0Gm));28+ T value; \
29- WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(dst0Gm), value);29+ if (Mode == 0) { \
30- } else {30+ value = READ_API<T>(reinterpret_cast<__gm__ T*>(src0Gm)); \
31- T value = ReadGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(src0Gm + sizeof(T)));31+ } else { \
32- WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(dst0Gm), value);32+ value = READ_API<T>(reinterpret_cast<__gm__ T*>(src0Gm + sizeof(T))); \
33- }33+ } \
34+ WRITE_API<T>(dstAddr, value); \
35+ } \
36+ \
37+ private: \
38+ GlobalTensor<T> src0Global; \
39+ GlobalTensor<T> dst0Global; \
40+ TPipe pipe; \
41+ TQue<QuePosition::VECIN, 1> inQueueX; \
42+ uint32_t dataSize = 0; \
43+ }; \
44+ template <typename T, uint8_t Mode> \
45+ void RunCase##TAG() \
46+ { \
47+ int srcByteSize = sizeof(T); \
48+ int dstByteSize = sizeof(T); \
49+ int dataSize = 256; \
50+ uint8_t dstGm0[dataSize * dstByteSize] = {0}; \
51+ uint8_t srcGm0[dataSize * srcByteSize] = {0}; \
52+ \
53+ KernelGmAccess##TAG<T, Mode> op; \
54+ op.Process(dstGm0, srcGm0, dataSize); \
34 }55 }
35 56 
36-private:57+DEFINE_GM_ACCESS_CASE(ReadGmBypassDCache, WriteGmBypassDCache, New)
37- GlobalTensor<T> src0Global;58+DEFINE_GM_ACCESS_CASE(ReadGmByPassDCache, WriteGmByPassDCache, Old)
38- GlobalTensor<T> dst0Global;
39- TPipe pipe;
40- TQue<QuePosition::VECIN, 1> inQueueX;
41- uint32_t dataSize = 0;
42-};
43 59 
44struct MicroGmAccessParams {60struct MicroGmAccessParams {
45 void (*CallFunc)();61 void (*CallFunc)();
46};62};
47 63 
48-template <typename T, uint8_t mode>
49-void RunCase()
50-{
51- int srcByteSize = sizeof(T);
52- int dstByteSize = sizeof(T);
53- int dataSize = 256;
54- uint8_t dstGm0[dataSize * dstByteSize] = {0};
55- uint8_t srcGm0[dataSize * srcByteSize] = {0};
56- 
57- KernelGmAccess<T, mode> op;
58- op.Process(dstGm0, srcGm0, dataSize);
59-}
60- 
61class MicroGmAccessTestSuite : public testing::Test, public testing::WithParamInterface<MicroGmAccessParams> {64class MicroGmAccessTestSuite : public testing::Test, public testing::WithParamInterface<MicroGmAccessParams> {
62protected:65protected:
63 void SetUp() {}66 void SetUp() {}
@@ -67,14 +70,22 @@ protected:
67INSTANTIATE_TEST_CASE_P(70INSTANTIATE_TEST_CASE_P(
68 MicroGmAccessTestCases, MicroGmAccessTestSuite,71 MicroGmAccessTestCases, MicroGmAccessTestSuite,
69 ::testing::Values(72 ::testing::Values(
70- MicroGmAccessParams{RunCase<uint64_t, 0>}, MicroGmAccessParams{RunCase<int64_t, 0>},73+ MicroGmAccessParams{RunCaseNew<uint64_t, 0>}, MicroGmAccessParams{RunCaseNew<int64_t, 0>},
71- MicroGmAccessParams{RunCase<uint32_t, 0>}, MicroGmAccessParams{RunCase<int32_t, 0>},74+ MicroGmAccessParams{RunCaseNew<uint32_t, 0>}, MicroGmAccessParams{RunCaseNew<int32_t, 0>},
72- MicroGmAccessParams{RunCase<uint16_t, 0>}, MicroGmAccessParams{RunCase<int16_t, 0>},75+ MicroGmAccessParams{RunCaseNew<uint16_t, 0>}, MicroGmAccessParams{RunCaseNew<int16_t, 0>},
73- MicroGmAccessParams{RunCase<uint8_t, 0>}, MicroGmAccessParams{RunCase<int8_t, 0>},76+ MicroGmAccessParams{RunCaseNew<uint8_t, 0>}, MicroGmAccessParams{RunCaseNew<int8_t, 0>},
74- MicroGmAccessParams{RunCase<uint64_t, 1>}, MicroGmAccessParams{RunCase<int64_t, 1>},77+ MicroGmAccessParams{RunCaseNew<uint64_t, 1>}, MicroGmAccessParams{RunCaseNew<int64_t, 1>},
75- MicroGmAccessParams{RunCase<uint32_t, 1>}, MicroGmAccessParams{RunCase<int32_t, 1>},78+ MicroGmAccessParams{RunCaseNew<uint32_t, 1>}, MicroGmAccessParams{RunCaseNew<int32_t, 1>},
76- MicroGmAccessParams{RunCase<uint16_t, 1>}, MicroGmAccessParams{RunCase<int16_t, 1>},79+ MicroGmAccessParams{RunCaseNew<uint16_t, 1>}, MicroGmAccessParams{RunCaseNew<int16_t, 1>},
77- MicroGmAccessParams{RunCase<uint8_t, 1>}, MicroGmAccessParams{RunCase<int8_t, 1>}));80+ MicroGmAccessParams{RunCaseNew<uint8_t, 1>}, MicroGmAccessParams{RunCaseNew<int8_t, 1>},
81+ MicroGmAccessParams{RunCaseOld<uint64_t, 0>}, MicroGmAccessParams{RunCaseOld<int64_t, 0>},
82+ MicroGmAccessParams{RunCaseOld<uint32_t, 0>}, MicroGmAccessParams{RunCaseOld<int32_t, 0>},
83+ MicroGmAccessParams{RunCaseOld<uint16_t, 0>}, MicroGmAccessParams{RunCaseOld<int16_t, 0>},
84+ MicroGmAccessParams{RunCaseOld<uint8_t, 0>}, MicroGmAccessParams{RunCaseOld<int8_t, 0>},
85+ MicroGmAccessParams{RunCaseOld<uint64_t, 1>}, MicroGmAccessParams{RunCaseOld<int64_t, 1>},
86+ MicroGmAccessParams{RunCaseOld<uint32_t, 1>}, MicroGmAccessParams{RunCaseOld<int32_t, 1>},
87+ MicroGmAccessParams{RunCaseOld<uint16_t, 1>}, MicroGmAccessParams{RunCaseOld<int16_t, 1>},
88+ MicroGmAccessParams{RunCaseOld<uint8_t, 1>}, MicroGmAccessParams{RunCaseOld<int8_t, 1>}));
78 89 
79TEST_P(MicroGmAccessTestSuite, MicroGmAccessTestCase)90TEST_P(MicroGmAccessTestSuite, MicroGmAccessTestCase)
80{91{
Mtests/api/basic_api/ascendc_case_mc62cm12aa/ascendc_case_mc62cm12aa_basic/test_global_memory_access.cpp+53-43
@@ -7,7 +7,6 @@
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10- 
11#include <gtest/gtest.h>10#include <gtest/gtest.h>
12#include <type_traits>11#include <type_traits>
13#include "kernel_operator.h"12#include "kernel_operator.h"
@@ -18,47 +17,50 @@ using AscendC::Reg::MaskReg;
18using AscendC::Reg::RegTensor;17using AscendC::Reg::RegTensor;
19using AscendC::Reg::UpdateMask;18using AscendC::Reg::UpdateMask;
20 19 
21-template <typename T, uint8_t Mode>20+#define DEFINE_GM_ACCESS_CASE(READ_API, WRITE_API, TAG) \
22-class KernelGmAccess {21+ template <typename T, uint8_t Mode> \
23-public:22+ class KernelGmAccess##TAG { \
24- __aicore__ inline KernelGmAccess() {}23+ public: \
25- 24+ __aicore__ inline KernelGmAccess##TAG() {} \
26- __aicore__ inline void Process(GM_ADDR dst0_gm, GM_ADDR src0_gm, uint32_t calcount)25+ __aicore__ inline void Process(GM_ADDR dst0Gm, GM_ADDR src0Gm, uint32_t calcount) \
27- {26+ { \
28- if (Mode == 0) {27+ __gm__ T* dstAddr = reinterpret_cast<__gm__ T*>(dst0Gm); \
29- T value = ReadGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(src0_gm));28+ T value; \
30- WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(dst0_gm), value);29+ if (Mode == 0) { \
31- } else {30+ value = READ_API<T>(reinterpret_cast<__gm__ T*>(src0Gm)); \
32- T value = ReadGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(src0_gm + sizeof(T)));31+ } else { \
33- WriteGmByPassDCache<T>(reinterpret_cast<__gm__ T*>(dst0_gm), value);32+ value = READ_API<T>(reinterpret_cast<__gm__ T*>(src0Gm + sizeof(T))); \
34- }33+ } \
34+ WRITE_API<T>(dstAddr, value); \
35+ } \
36+ \
37+ private: \
38+ GlobalTensor<T> src0Global; \
39+ GlobalTensor<T> dst0Global; \
40+ TPipe pipe; \
41+ TQue<QuePosition::VECIN, 1> inQueueX; \
42+ uint32_t dataSize = 0; \
43+ }; \
44+ template <typename T, uint8_t Mode> \
45+ void RunCase##TAG() \
46+ { \
47+ int srcByteSize = sizeof(T); \
48+ int dstByteSize = sizeof(T); \
49+ int dataSize = 256; \
50+ uint8_t dstGm0[dataSize * dstByteSize] = {0}; \
51+ uint8_t srcGm0[dataSize * srcByteSize] = {0}; \
52+ \
53+ KernelGmAccess##TAG<T, Mode> op; \
54+ op.Process(dstGm0, srcGm0, dataSize); \
35 }55 }
36 56 
37-private:57+DEFINE_GM_ACCESS_CASE(ReadGmBypassDCache, WriteGmBypassDCache, New)
38- GlobalTensor<T> src0_global;58+DEFINE_GM_ACCESS_CASE(ReadGmByPassDCache, WriteGmByPassDCache, Old)
39- GlobalTensor<T> dst0_global;
40- TPipe pipe;
41- TQue<QuePosition::VECIN, 1> inQueueX;
42- uint32_t dataSize = 0;
43-};
44 59 
45struct MicroGmAccessParams {60struct MicroGmAccessParams {
46 void (*CallFunc)();61 void (*CallFunc)();
47};62};
48 63 
49-template <typename T, uint8_t Mode>
50-void RunCase()
51-{
52- int src_byte_size = sizeof(T);
53- int dst_byte_size = sizeof(T);
54- int data_size = 256;
55- uint8_t dstGm0[data_size * dst_byte_size] = {0};
56- uint8_t srcGm0[data_size * src_byte_size] = {0};
57- 
58- KernelGmAccess<T, Mode> op;
59- op.Process(dstGm0, srcGm0, data_size);
60-}
61- 
62class MicroGmAccessTestSuite : public testing::Test, public testing::WithParamInterface<MicroGmAccessParams> {64class MicroGmAccessTestSuite : public testing::Test, public testing::WithParamInterface<MicroGmAccessParams> {
63protected:65protected:
64 void SetUp() {}66 void SetUp() {}
@@ -68,14 +70,22 @@ protected:
68INSTANTIATE_TEST_CASE_P(70INSTANTIATE_TEST_CASE_P(
69 MicroGmAccessTestCases, MicroGmAccessTestSuite,71 MicroGmAccessTestCases, MicroGmAccessTestSuite,
70 ::testing::Values(72 ::testing::Values(
71- MicroGmAccessParams{RunCase<uint64_t, 0>}, MicroGmAccessParams{RunCase<int64_t, 0>},73+ MicroGmAccessParams{RunCaseNew<uint64_t, 0>}, MicroGmAccessParams{RunCaseNew<int64_t, 0>},
72- MicroGmAccessParams{RunCase<uint32_t, 0>}, MicroGmAccessParams{RunCase<int32_t, 0>},74+ MicroGmAccessParams{RunCaseNew<uint32_t, 0>}, MicroGmAccessParams{RunCaseNew<int32_t, 0>},
73- MicroGmAccessParams{RunCase<uint16_t, 0>}, MicroGmAccessParams{RunCase<int16_t, 0>},75+ MicroGmAccessParams{RunCaseNew<uint16_t, 0>}, MicroGmAccessParams{RunCaseNew<int16_t, 0>},
74- MicroGmAccessParams{RunCase<uint8_t, 0>}, MicroGmAccessParams{RunCase<int8_t, 0>},76+ MicroGmAccessParams{RunCaseNew<uint8_t, 0>}, MicroGmAccessParams{RunCaseNew<int8_t, 0>},
75- MicroGmAccessParams{RunCase<uint64_t, 1>}, MicroGmAccessParams{RunCase<int64_t, 1>},77+ MicroGmAccessParams{RunCaseNew<uint64_t, 1>}, MicroGmAccessParams{RunCaseNew<int64_t, 1>},
76- MicroGmAccessParams{RunCase<uint32_t, 1>}, MicroGmAccessParams{RunCase<int32_t, 1>},78+ MicroGmAccessParams{RunCaseNew<uint32_t, 1>}, MicroGmAccessParams{RunCaseNew<int32_t, 1>},
77- MicroGmAccessParams{RunCase<uint16_t, 1>}, MicroGmAccessParams{RunCase<int16_t, 1>},79+ MicroGmAccessParams{RunCaseNew<uint16_t, 1>}, MicroGmAccessParams{RunCaseNew<int16_t, 1>},
78- MicroGmAccessParams{RunCase<uint8_t, 1>}, MicroGmAccessParams{RunCase<int8_t, 1>}));80+ MicroGmAccessParams{RunCaseNew<uint8_t, 1>}, MicroGmAccessParams{RunCaseNew<int8_t, 1>},
81+ MicroGmAccessParams{RunCaseOld<uint64_t, 0>}, MicroGmAccessParams{RunCaseOld<int64_t, 0>},
82+ MicroGmAccessParams{RunCaseOld<uint32_t, 0>}, MicroGmAccessParams{RunCaseOld<int32_t, 0>},
83+ MicroGmAccessParams{RunCaseOld<uint16_t, 0>}, MicroGmAccessParams{RunCaseOld<int16_t, 0>},
84+ MicroGmAccessParams{RunCaseOld<uint8_t, 0>}, MicroGmAccessParams{RunCaseOld<int8_t, 0>},
85+ MicroGmAccessParams{RunCaseOld<uint64_t, 1>}, MicroGmAccessParams{RunCaseOld<int64_t, 1>},
86+ MicroGmAccessParams{RunCaseOld<uint32_t, 1>}, MicroGmAccessParams{RunCaseOld<int32_t, 1>},
87+ MicroGmAccessParams{RunCaseOld<uint16_t, 1>}, MicroGmAccessParams{RunCaseOld<int16_t, 1>},
88+ MicroGmAccessParams{RunCaseOld<uint8_t, 1>}, MicroGmAccessParams{RunCaseOld<int8_t, 1>}));
79 89 
80TEST_P(MicroGmAccessTestSuite, MicroGmAccessTestCase)90TEST_P(MicroGmAccessTestSuite, MicroGmAccessTestCase)
81{91{
Mtests/api/basic_api/ascendc_header_checker/kernel_operator_scalar_intf.cpp+3-3
@@ -26,11 +26,11 @@ extern "C" __global__ __aicore__ void KernelTestGetUintDivMagicAndShift1()
26#endif26#endif
27 27 
28#if __NPU_ARCH__ == 2201 || (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102)28#if __NPU_ARCH__ == 2201 || (__NPU_ARCH__ == 3510) || (__NPU_ARCH__ == 5102)
29-// __aicore__ inline void WriteGmByPassDCache(__gm__ T* addr, T value);29+// __aicore__ inline void WriteGmBypassDCache(__gm__ T* addr, T value);
30-extern "C" __global__ __aicore__ void KernelTestWriteGmByPassDCache1()30+extern "C" __global__ __aicore__ void KernelTestWriteGmBypassDCache1()
31{31{
32 __gm__ int32_t* addr;32 __gm__ int32_t* addr;
33 int32_t value = 0;33 int32_t value = 0;
34- AscendC::WriteGmByPassDCache(addr, value);34+ AscendC::WriteGmBypassDCache(addr, value);
35}35}
36#endif36#endif
Mtests/api/basic_api/ascendc_host_header_checker/host_header_checker.cpp+2-2
@@ -125,8 +125,8 @@ static void test_host_kernel_operator_scalar_intf()
125{125{
126 using AscendC::Cast;126 using AscendC::Cast;
127 using AscendC::GetUintDivMagicAndShift;127 using AscendC::GetUintDivMagicAndShift;
128- using AscendC::ReadGmByPassDCache;128+ using AscendC::ReadGmBypassDCache;
129- using AscendC::WriteGmByPassDCache;129+ using AscendC::WriteGmBypassDCache;
130}130}
131 131 
132static void test_host_kernel_operator_sys_var_intf()132static void test_host_kernel_operator_sys_var_intf()
Mtests/api/basic_api/ascendc_host_header_checker/host_header_kernal_operator_intf_checker.cpp+2-2
@@ -176,12 +176,12 @@ static void test_host_kernel_operator_scalar_intf()
176 using AscendC::GetBitCount;176 using AscendC::GetBitCount;
177 using AscendC::GetSFFValue;177 using AscendC::GetSFFValue;
178 using AscendC::GetUintDivMagicAndShift;178 using AscendC::GetUintDivMagicAndShift;
179- using AscendC::ReadGmByPassDCache;179+ using AscendC::ReadGmBypassDCache;
180 using AscendC::ScalarCast;180 using AscendC::ScalarCast;
181 using AscendC::ScalarCountLeadingZero;181 using AscendC::ScalarCountLeadingZero;
182 using AscendC::ScalarGetCountOfValue;182 using AscendC::ScalarGetCountOfValue;
183 using AscendC::ScalarGetSFFValue;183 using AscendC::ScalarGetSFFValue;
184- using AscendC::WriteGmByPassDCache;184+ using AscendC::WriteGmBypassDCache;
185}185}
186 186 
187static void test_host_kernel_operator_vec_reduce_intf()187static void test_host_kernel_operator_vec_reduce_intf()