docs: consolidate assert API documentation-asc-devkit-AtomGit

Mdocs/zh/api/README.md+1-3

	@@ -663,8 +663,6 @@
663	- [PrintTimeStamp](SIMD-API/basic_api/debug_interface/onboard_print/PrintTimeStamp.md)	663	- [PrintTimeStamp](SIMD-API/basic_api/debug_interface/onboard_print/PrintTimeStamp.md)
664		664
665	- [异常检测](SIMD-API/basic_api/debug_interface/exception_detection/exception_detection.md)	665	- [异常检测](SIMD-API/basic_api/debug_interface/exception_detection/exception_detection.md)
666	- - [ascendc\_assert](SIMD-API/basic_api/debug_interface/exception_detection/ascendc_assert.md)
667	- - [assert](SIMD-API/basic_api/debug_interface/exception_detection/assert.md)
668	- [Trap](SIMD-API/basic_api/debug_interface/exception_detection/Trap.md)	666	- [Trap](SIMD-API/basic_api/debug_interface/exception_detection/Trap.md)
669	- [CheckLocalMemoryIA\(ISASI\)](SIMD-API/basic_api/debug_interface/exception_detection/CheckLocalMemoryIA_ISASI.md)	667	- [CheckLocalMemoryIA\(ISASI\)](SIMD-API/basic_api/debug_interface/exception_detection/CheckLocalMemoryIA_ISASI.md)
670		668
	@@ -3194,7 +3192,7 @@
3194		3192
3195	- [调测接口](Utils-API/tuning_interface/tuning_interface.md)	3193	- [调测接口](Utils-API/tuning_interface/tuning_interface.md)
3196	- [printf](Utils-API/tuning_interface/printf.md)	3194	- [printf](Utils-API/tuning_interface/printf.md)
3197	- - [assert](Utils-API/tuning_interface/assert.md)	3195	+ - [assert/ascendc\_assert](Utils-API/tuning_interface/assert.md)
3198	- [\_\_trap](Utils-API/tuning_interface/__trap.md)	3196	- [\_\_trap](Utils-API/tuning_interface/__trap.md)
3199	- [asc\_dump](Utils-API/tuning_interface/asc_dump.md)	3197	- [asc\_dump](Utils-API/tuning_interface/asc_dump.md)
3200	- [clock](Utils-API/tuning_interface/clock.md)	3198	- [clock](Utils-API/tuning_interface/clock.md)

Mdocs/zh/api/SIMD-API/adv_api/HCCL_communication/HCCL_Kernel/HCCL_usage.md+1-1

	@@ -249,4 +249,4 @@ extern "C" __global__ __aicore__ void reduce_scatter_custom(
249	\| --- \| --- \|	249	\| --- \| --- \|
250	\| MC2_BUFFER_LOCATION \| 预留参数。计算和通信中间结果的Buffer存放位置。用户在Tiling侧可设置该字段。 \|	250	\| MC2_BUFFER_LOCATION \| 预留参数。计算和通信中间结果的Buffer存放位置。用户在Tiling侧可设置该字段。 \|
251		251
252	-提示：调试含有HCCL高阶API的算子时，在算子编译工程中增加编译选项-DASCENDC\_DEBUG，可以开启异常场景拦截的能力，具体内容请参考并使用[assert接口](../../../basic_api/debug_interface/exception_detection/assert.md)。	252	+提示：调试含有HCCL高阶API的算子时，在算子编译工程中增加编译选项-DASCENDC\_DEBUG，可以开启异常场景拦截的能力，具体内容请参考并使用[assert接口](../../../../Utils-API/tuning_interface/assert.md)。

Mdocs/zh/api/SIMD-API/basic_api/basic_api_list.md+0-2

	@@ -511,8 +511,6 @@
511	\| [DumpTensor](debug_interface/onboard_print/DumpTensor.md) \| 该接口可以打印Tensor的内容，同时支持打印自定义的标签（仅支持uint32_t数据类型的信息），比如打印当前行号等。 \|	511	\| [DumpTensor](debug_interface/onboard_print/DumpTensor.md) \| 该接口可以打印Tensor的内容，同时支持打印自定义的标签（仅支持uint32_t数据类型的信息），比如打印当前行号等。 \|
512	\| [DumpAccChkPoint](debug_interface/onboard_print/DumpAccChkPoint.md) \| 该接口Dump指定Tensor的内容。同时支持打印自定义的标签（仅支持uint32_t数据类型的信息），比如打印当前行号等。区别于DumpTensor，使用该接口可以支持指定偏移位置的Tensor打印。 \|	512	\| [DumpAccChkPoint](debug_interface/onboard_print/DumpAccChkPoint.md) \| 该接口Dump指定Tensor的内容。同时支持打印自定义的标签（仅支持uint32_t数据类型的信息），比如打印当前行号等。区别于DumpTensor，使用该接口可以支持指定偏移位置的Tensor打印。 \|
513	\| [PrintTimeStamp](debug_interface/onboard_print/PrintTimeStamp.md) \| 提供时间戳打点功能，用于在算子核函数（Kernel）代码中标记关键执行点。 \|	513	\| [PrintTimeStamp](debug_interface/onboard_print/PrintTimeStamp.md) \| 提供时间戳打点功能，用于在算子核函数（Kernel）代码中标记关键执行点。 \|
514	-\| [ascendc_assert](debug_interface/exception_detection/ascendc_assert.md) \| ascendc_assert提供了一种在CPU/NPU域实现断言功能的接口。当断言条件不满足时，系统会输出断言信息并格式化打印在屏幕上。 \|
515	-\| [assert](debug_interface/exception_detection/assert.md) \| 该接口实现CPU/NPU域assert断言功能。算子执行中，如果assert内部条件判断不为真，则输出assert条件并将输入的信息格式化打印在屏幕上。 \|
516	\| [Trap](debug_interface/exception_detection/Trap.md) \| 在核函数（Kernel）侧调用，NPU模式下会中断AI Core的运行，CPU模式下等同于assert。可用于核函数（Kernel）侧异常场景的调试。 \|	514	\| [Trap](debug_interface/exception_detection/Trap.md) \| 在核函数（Kernel）侧调用，NPU模式下会中断AI Core的运行，CPU模式下等同于assert。可用于核函数（Kernel）侧异常场景的调试。 \|
517	\| [CheckLocalMemoryIA(ISASI)](debug_interface/exception_detection/CheckLocalMemoryIA_ISASI.md) \| check设定范围内的UB读写行为，如果有设定范围的读写行为则会出现EXCEPTION报错，无设定范围的读写行为则不会报错。 \|	515	\| [CheckLocalMemoryIA(ISASI)](debug_interface/exception_detection/CheckLocalMemoryIA_ISASI.md) \| check设定范围内的UB读写行为，如果有设定范围的读写行为则会出现EXCEPTION报错，无设定范围的读写行为则不会报错。 \|
518	\| [GmAlloc](debug_interface/cpu_twin_debug/GmAlloc.md) \| 进行核函数（Kernel）的CPU侧运行验证时，用于创建共享内存：在/tmp目录下创建一个共享文件，并返回该文件的映射指针。 \|	516	\| [GmAlloc](debug_interface/cpu_twin_debug/GmAlloc.md) \| 进行核函数（Kernel）的CPU侧运行验证时，用于创建共享内存：在/tmp目录下创建一个共享文件，并返回该文件的映射指针。 \|

Ddocs/zh/api/SIMD-API/basic_api/debug_interface/exception_detection/ascendc_assert.md+0-84

	@@ -1,84 +0,0 @@
1	-# ascendc\_assert<a name="ZH-CN_TOPIC_0000002506195843"></a>
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	-<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/debug_interface/exception_detection/ascendc_assert_res.md#id1 -->
27	-
28	-## 功能说明<a name="section259105813316"></a>
29	-
30	-头文件路径为：`"utils/debug/asc_assert.h"`。
31	-
32	-ascendc_assert提供了一种在CPU/NPU域实现断言功能的接口。当断言条件不满足时，系统会输出断言信息并格式化打印在屏幕上。
33	-
34	-在算子核函数（Kernel）侧实现代码中需要增加断言的地方使用ascendc_assert检查代码，并格式化输出一些调测信息。
35	-
36	-> [!CAUTION]注意
37	-> 该接口主要用于调试分析，开启后会对算子性能产生一定影响，通常在调试阶段使用，生产环境建议关闭。<br>
38	-> 默认情况下，调用该接口就会打印相关内容，开发者可以参考[关闭ASCENDC_DUMP说明](../disable_ascendc_dump.md)，按需关闭该接口功能。
39	-
40	-## 函数原型<a name="section2067518173415"></a>
41	-
42	-```cpp
43	-#define ascendc_assert(expr)
44	-#define ascendc_assert(expr, __gm__ const char *fmt, Args&&... args)
45	-```
46	-
47	-## 参数说明<a name="section158061867342"></a>
48	-
49	-表1 参数说明
50	-
51	-\| 参数名称 \| 输入/输出 \| 描述 \|
52	-\| ------ \| ------ \| ------ \|
53	-\| expr \| 输入 \| assert断言是否终止程序的条件。为true则程序继续执行，为false则终止程序。 \|
54	-\| fmt \| 输入 \| 格式控制字符串，包含两种类型的对象：普通字符和转换说明。<br>•普通字符将原样不动地打印输出。<br>•转换说明并不直接输出而是用于控制printf中参数的转换和打印。每个转换说明都由一个百分号字符（%）开始，以转换说明结束，从而说明输出数据的类型。<br>•支持的转换类型包括：<br>&ensp;&ensp;▪%d / %i：输出十进制数，支持打印的数据类型：bool、int8_t、int16_t、int32_t、int64_t。<br>&ensp;&ensp;▪%f：输出实数，支持打印的数据类型：float、half。<br>&ensp;&ensp;▪%x：输出十六进制整数，支持打印的数据类型：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。<br>&ensp;&ensp;▪%s：输出字符串。<br>&ensp;&ensp;▪%u：输出unsigned类型数据，支持打印的数据类型：bool、uint8_t、uint16_t、uint32_t、uint64_t。<br>&ensp;&ensp;▪%p：输出指针地址。<br>注意：<br>•上文列出的数据类型是NPU域调试支持的数据类型，CPU域调试时，支持的数据类型和C/C++规范保持一致。<br>•在转换类型为%x，即输出十六进制整数时，NPU域上的输出为64位，CPU域上的输出为32位。 \|
55	-\| args \| 输入 \| 附加参数，个数和类型可变的参数列表：根据不同的fmt字符串，函数可能需要一系列的附加参数，每个参数包含了一个要被插入的值，替换了fmt参数中指定的每个%标签。参数的个数应与%标签的个数相同。 \|
56	-
57	-## 数据类型
58	-
59	-无
60	-
61	-## 返回值说明<a name="section640mcpsimp"></a>
62	-
63	-无
64	-
65	-## 约束说明<a name="section43265506459"></a>
66	-
67	-单次调用本接口打印的数据总量不可超过30KB。使用时应注意，如果超出这个限制，则数据不会被打印。
68	-
69	-## 调用示例<a name="section82241477610"></a>
70	-
71	-接口调用示例
72	-
73	-```cpp
74	-int assertFlag = 10;
75	-
76	-// 格式化打印
77	-ascendc_assert(assertFlag != 10, "The assertFlag value is %d.\n", assertFlag);
78	-```
79	-
80	-程序运行时会触发assert，打印效果如下：
81	-
82	-```plain
83	-[ASSERT] /home/.../add_custom.cpp:44: Assertion `assertFlag != 10' The assertFlag value is 10.
84	-```

Ddocs/zh/api/SIMD-API/basic_api/debug_interface/exception_detection/assert.md+0-94

	@@ -1,94 +0,0 @@
1	-# assert<a name="ZH-CN_TOPIC_0000001920717221"></a>
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	-<!-- npu="x90" id8 -->
27	-- Kirin X90：支持
28	-<!-- end id8 -->
29	-<!-- npu="9030" id9 -->
30	-- Kirin 9030：支持
31	-<!-- end id9 -->
32	-<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/debug_interface/exception_detection/assert_res.md#id1 -->
33	-
34	-## 功能说明<a name="section259105813316"></a>
35	-
36	-头文件路径为：`"utils/debug/asc_assert.h"`。
37	-
38	-该接口实现CPU/NPU域assert断言功能。算子执行中，如果assert内部条件判断不为真，则输出assert条件并将输入的信息格式化打印在屏幕上。
39	-
40	-在算子核函数（Kernel）侧实现代码中需要增加断言的地方使用assert检查代码，并格式化输出一些调测信息。示例如下：
41	-
42	-```cpp
43	-int assertFlag = 10;
44	-
45	-assert(assertFlag == 10);
46	-```
47	-
48	-> [!CAUTION]注意
49	-> 该接口主要用于调试分析，开启后会对算子性能产生一定影响，通常在调试阶段使用，生产环境建议关闭。<br>
50	-> 默认情况下，调用该接口就会打印相关内容，开发者可以参考[关闭ASCENDC_DUMP说明](../disable_ascendc_dump.md)，按需关闭该接口功能。
51	-
52	-## 函数原型<a name="section2067518173415"></a>
53	-
54	-```cpp
55	-define assert(expr)
56	-define assert(expr, __gm__ const char *fmt, Args&&... args)
57	-```
58	-
59	-## 参数说明<a name="section158061867342"></a>
60	-
61	-表1 参数说明
62	-
63	-\| 参数名称 \| 输入/输出 \| 描述 \|
64	-\| ------ \| ------ \| ------ \|
65	-\| expr \| 输入 \| assert断言是否终止程序的条件。为true则程序继续执行，为false则终止程序。 \|
66	-\| fmt \| 输入 \| 格式控制字符串，包含两种类型的对象：普通字符和转换说明。<br>•普通字符将原样不动地打印输出。<br>•转换说明并不直接输出而是用于控制printf中参数的转换和打印。每个转换说明都由一个百分号字符（%）开始，以转换说明结束，从而说明输出数据的类型。<br>•支持的转换类型包括：<br>&ensp;&ensp;▪%d / %i：输出十进制数，支持打印的数据类型：bool、int8_t、int16_t、int32_t、int64_t。<br>&ensp;&ensp;▪%f：输出实数，支持打印的数据类型：float、half。<br>&ensp;&ensp;▪%x：输出十六进制整数，支持打印的数据类型：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。<br>&ensp;&ensp;▪%s：输出字符串。<br>&ensp;&ensp;▪%u：输出unsigned类型数据，支持打印的数据类型：bool、uint8_t、uint16_t、uint32_t、uint64_t。<br>&ensp;&ensp;▪%p：输出指针地址。<br>注意：<br>•上文列出的数据类型是NPU域调试支持的数据类型，CPU域调试时，支持的数据类型和C/C++规范保持一致。<br>•在转换类型为%x，即输出十六进制整数时，NPU域上的输出为64位，CPU域上的输出为32位。 \|
67	-\| args \| 输入 \| 附加参数，个数和类型可变的参数列表：根据不同的fmt字符串，函数可能需要一系列的附加参数，每个参数包含了一个要被插入的值，替换了fmt参数中指定的每个%标签。参数的个数应与%标签的个数相同。 \|
68	-
69	-## 数据类型
70	-
71	-无
72	-
73	-## 返回值说明<a name="section640mcpsimp"></a>
74	-
75	-无
76	-
77	-## 约束说明<a name="section43265506459"></a>
78	-
79	-- 单次调用本接口打印的数据总量不可超过30KB。使用时应注意，如果超出这个限制，则数据不会被打印。
80	-- 在cmake工程的release模式下，默认使能NDEBUG宏定义，assert打印信息失效。
81	-
82	-## 调用示例<a name="section82241477610"></a>
83	-
84	-```cpp
85	-int assertFlag = 10;
86	-// 断言条件
87	-assert(assertFlag != 10);
88	-```
89	-
90	-程序运行时会触发assert，打印效果如下：
91	-
92	-```plain
93	-[ASSERT] /home/.../add_custom.cpp:44: Assertion `assertFlag != 10'
94	-```

Mdocs/zh/api/SIMD-API/basic_api/debug_interface/exception_detection/exception_detection.md+0-4

	@@ -1,9 +1,5 @@
1	# 异常检测<a name="ZH-CN_TOPIC_0000002512364579"></a>	1	# 异常检测<a name="ZH-CN_TOPIC_0000002512364579"></a>
2		2
3	-- [ascendc\_assert](ascendc_assert.md)
4	-
5	-- [assert](assert.md)
6	-
7	- [Trap](Trap.md)	3	- [Trap](Trap.md)
8		4
9	- [CheckLocalMemoryIA\(ISASI\)](CheckLocalMemoryIA_ISASI.md)	5	- [CheckLocalMemoryIA\(ISASI\)](CheckLocalMemoryIA_ISASI.md)

Mdocs/zh/api/Utils-API/tuning_interface/asc_dump.md+2-1

	@@ -145,8 +145,9 @@ __simd_callee__ inline void asc_dump(__ubuf__ T* input, uint32_t desc, uint32_t
145	- SIMD场景下，单次调用本接口打印的数据总量不可超过打印大小限制，默认为30KB。使用时应注意，如果超出这个限制，则数据不会被打印。您可以通过acl.json中的`"simd_printf_fifo_size_per_core"`字段进行配置，配置范围最小为1KB，最大为64MB（可通过[aclInit](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/latest/API/runtimeapi/aclcppdevg_03_0022.html)接口调整）。当打印数据量较大时，建议增加缓存空间。pytorch调用和算子入图场景暂不支持该配置。	145	- SIMD场景下，单次调用本接口打印的数据总量不可超过打印大小限制，默认为30KB。使用时应注意，如果超出这个限制，则数据不会被打印。您可以通过acl.json中的`"simd_printf_fifo_size_per_core"`字段进行配置，配置范围最小为1KB，最大为64MB（可通过[aclInit](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/latest/API/runtimeapi/aclcppdevg_03_0022.html)接口调整）。当打印数据量较大时，建议增加缓存空间。pytorch调用和算子入图场景暂不支持该配置。
146		146
147	<!-- npu="950" id16 -->	147	<!-- npu="950" id16 -->
148	-- 在`simd_vf`场景下，每个AIV核使用2KB预留UB空间作为FIFO临时缓冲区。同一次`asc_vf_call`中，`printf`、`asc_dump`和`assert`产生的累计调测数据可以超过2KB，但每条完整编码数据必须能放入该缓冲区；否则，该条数据不会打印。	148	+- SIMD VF场景下，每个AIV核在单次`asc_vf_call`执行期间使用2KB预留UB空间临时保存调测数据。同一次`asc_vf_call`中的`assert`、`ascendc_assert`、`printf`和`asc_dump`共享该空间。该空间中的数据传输完成后会被复用，因此上述接口产生的累计调测数据可以超过2KB。单条调测数据必须能完整保存在该空间中，否则该条数据不会打印。
149	- 每次调用`simd_vf`的`asc_dump`时，除实际dump数据外，还会固定占用72字节的管理信息；实际dump数据需要按32字节向上对齐。	149	- 每次调用`simd_vf`的`asc_dump`时，除实际dump数据外，还会固定占用72字节的管理信息；实际dump数据需要按32字节向上对齐。
		150	+- SIMD VF场景下，`simd_printf_fifo_size_per_core`建议配置为3KB以上。配置过小且打印数据量较大时，部分调测数据不会被打印。
150	<!-- end id16 -->	151	<!-- end id16 -->
151		152
152	## 调用示例	153	## 调用示例

Mdocs/zh/api/Utils-API/tuning_interface/assert.md+140-56

	@@ -1,4 +1,4 @@
1	-# assert	1	+# assert/ascendc_assert
2		2
3	## 产品支持情况	3	## 产品支持情况
4		4
	@@ -15,50 +15,104 @@
15	- Atlas 200I/500 A2 推理产品：不支持	15	- Atlas 200I/500 A2 推理产品：不支持
16	<!-- end id4 -->	16	<!-- end id4 -->
17	<!-- npu="310p" id5 -->	17	<!-- npu="310p" id5 -->
18	-- Atlas 推理系列产品AI Core：不支持	18	+- Atlas 推理系列产品AI Core：支持
19	-- Atlas 推理系列产品Vector Core：不支持
20	<!-- end id5 -->	19	<!-- end id5 -->
21	-<!-- npu="910" id6 -->	20	+<!-- npu="310p" id6 -->
22	-- Atlas 训练系列产品：不支持	21	+- Atlas 推理系列产品Vector Core：支持
		Y Y YANXI_ZHAO15 天前 vector core确定支持么？建议你找一个用例进行验证，我记得是没有做支持的。
23	<!-- end id6 -->	22	<!-- end id6 -->
24	-<!-- npu="x90" id7 -->	23	+<!-- npu="910" id7 -->
25	-- Kirin X90：不支持	24	+- Atlas 训练系列产品：不支持
26	<!-- end id7 -->	25	<!-- end id7 -->
27	-<!-- npu="9030" id8 -->	26	+<!-- npu="x90" id8 -->
28	-- Kirin 9030：不支持	27	+- Kirin X90：不支持
29	<!-- end id8 -->	28	<!-- end id8 -->
		29	+<!-- npu="9030" id9 -->
		30	+- Kirin 9030：不支持
		31	+<!-- end id9 -->
		32	+
		33	+不同产品对各编程场景的支持情况如下表所示。
		34	+
		35	+表1 支持能力
		36	+
		37	+<table>
		38	+ <tr>
		39	+ <th>产品</th>
		40	+ <th>SIMD</th>
		41	+ <th>SIMT VF</th>
		42	+ <th>SIMD VF</th>
		43	+ </tr>
		44	+ <!-- npu="950" id10 -->
		45	+ <tr>
		46	+ <td>Ascend 950PR/Ascend 950DT</td>
		47	+ <td>支持</td>
		48	+ <td>支持</td>
		49	+ <td>支持</td>
		50	+ </tr>
		51	+ <!-- end id10 -->
		52	+ <!-- npu="A3" id11 -->
		53	+ <tr>
		54	+ <td>Atlas A3 训练系列产品/Atlas A3 推理系列产品</td>
		55	+ <td>支持</td>
		56	+ <td>不支持</td>
		57	+ <td>不支持</td>
		58	+ </tr>
		59	+ <!-- end id11 -->
		60	+ <!-- npu="910b" id12 -->
		61	+ <tr>
		62	+ <td>Atlas A2 训练系列产品/Atlas A2 推理系列产品</td>
		63	+ <td>支持</td>
		64	+ <td>不支持</td>
		65	+ <td>不支持</td>
		66	+ </tr>
		67	+ <!-- end id12 -->
		68	+ <!-- npu="310p" id13 -->
		69	+ <tr>
		70	+ <td>Atlas 推理系列产品 AI Core</td>
		71	+ <td>支持</td>
		72	+ <td>不支持</td>
		73	+ <td>不支持</td>
		74	+ </tr>
		75	+ <!-- end id13 -->
		76	+ <!-- npu="310p" id14 -->
		77	+ <tr>
		78	+ <td>Atlas 推理系列产品 Vector Core</td>
		79	+ <td>支持</td>
		80	+ <td>不支持</td>
		81	+ <td>不支持</td>
		82	+ </tr>
		83	+ <!-- end id14 -->
		84	+</table>
30		85
31	## 功能说明	86	## 功能说明
32		87
33	-本接口在SIMD和SIMT调试场景下提供assert断言功能。在算子核函数（Kernel）侧的实现代码中，如果assert的内部条件判断不为真，则会输出assert条件，并将输入的信息格式化打印在屏幕上，同时算子运行失败。	88	+头文件路径为：`"utils/debug/asc_assert.h"`。
34		89
35	-在算子核函数（Kernel）侧代码的适当位置使用assert进行断言检查，并格式化输出一些调试信息。示例如下：	90	+`assert`和`ascendc_assert`用于在Device侧代码中进行断言检查。断言条件成立时，后续代码继续执行；断言条件不成立时，接口打印断言失败信息并触发异常，算子执行失败。
36		91
37	-```cpp	92	+需要在断言失败时输出自定义错误信息，可以在条件后传入格式字符串和对应参数。`assert`和`ascendc_assert`提供相同的断言能力，可以任选其一使用。
38	-int assertFlag = 10;
39		93
40	-assert(assertFlag != 10);	94	+> [!CAUTION]注意
41	-```	95	+> 该接口用于调测，使用时会影响算子性能，生产环境建议通过设置[ASCENDC_DUMP=0](../../SIMD-API/basic_api/debug_interface/disable_ascendc_dump.md)关闭。
42	-
43	-打印信息示例如下：
44	-
45	-```
46	-[ASSERT] /home/.../add_custom.cpp:44: : Assertion `assertFlag != 10' failed.
47	-```
48	-
49	-请注意，assert接口的打印功能对算子的实际运行性能有影响。
50		96
51	## 函数原型	97	## 函数原型
52		98
53	```cpp	99	```cpp
54	assert(expr)	100	assert(expr)
		101	+assert(expr, fmt, args...)
		102	+
		103	+ascendc_assert(expr)
		104	+ascendc_assert(expr, fmt, args...)
55	```	105	```
56		106
57	## 参数说明	107	## 参数说明
58		108
		109	+表2 参数说明
		110	+
59	\| 参数名 \| 输入/输出 \| 描述 \|	111	\| 参数名 \| 输入/输出 \| 描述 \|
60	\| --- \| --- \| --- \|	112	\| --- \| --- \| --- \|
61	-\| expr \| 输入 \| assert断言是否终止程序的条件。条件为true则程序继续执行，条件为false则终止程序。 \|	113	+\| expr \| 输入 \| 断言条件。为`true`时，Kernel继续运行；为`false`时，接口打印断言失败信息并触发异常。 \|
		114	+\| fmt \| 输入 \| 可选。自定义错误信息的格式控制字符串，包含普通字符和转换说明。普通字符原样输出；转换说明以百分号（%）开始，用于控制`args`中对应参数的转换和输出。支持的转换类型和数据类型参见[printf](printf.md)的参数说明。 \|
		115	+\| args \| 输入 \| 可选。与`fmt`中的转换说明对应的附加参数。每个参数替换`fmt`中对应的转换说明，参数数量和类型应与转换说明匹配。 \|
62		116
63	## 返回值说明	117	## 返回值说明
64		118
	@@ -66,47 +120,59 @@ assert(expr)
66		120
67	## 约束说明	121	## 约束说明
68		122
69	-- 该接口当前仅支持融合编译场景。	123	+### 通用约束
70	-- SIMD场景，如果开发者需要包含标准库头文件<cassert\>，请在"utils/debug/asc\_assert.h"头文件之前包含，避免assert符号冲突。
71	-<!-- npu="950,A3,910b" id13 -->
72	-- SIMT场景的支持情况如下：
73	- <!-- npu="950" id10 -->
74	- - Ascend 950PR/Ascend 950DT：支持。
75	- <!-- end id10 -->
76	- <!-- npu="A3" id11 -->
77	- - Atlas A3 训练系列产品/Atlas A3 推理系列产品：不支持。
78	- <!-- end id11 -->
79	- <!-- npu="910b" id12 -->
80	- - Atlas A2 训练系列产品/Atlas A2 推理系列产品：不支持。
81	- <!-- end id12 -->
82	-<!-- end id13 -->
83		124
84	-## 需要包含的头文件	125	+- 在CMake工程的Release模式下，默认使能`NDEBUG`宏定义，本接口不生效。
		M M munanhw9 天前能否不体现默认使能 NDEBUG定义
		126	+- 使用`assert`时，如需包含标准库头文件`<cassert>`，请先包含`<cassert>`，再包含`"kernel_operator.h"`或`"utils/debug/asc_assert.h"`，避免标准库覆盖Ascend C的`assert`定义。
85		127
86	-使用该接口需要包含"utils/debug/asc\_assert.h"头文件。	128	+### SIMD编程场景
		129	+
		130	+- 单次调用本接口打印的数据总量不可超过打印大小限制，默认为30KB。超出限制时，断言失败信息不会打印，但接口仍会触发异常。可以通过[aclInit接口](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/latest/API/runtimeapi/aclcppdevg_03_0022.html)中的`simd_printf_fifo_size_per_core`字段配置，配置范围为`[1KB,64MB]`。PyTorch调用和算子入图场景暂不支持该配置。
		131	+
		132	+### SIMT VF编程场景
		133	+
		134	+不输出`fmt`和`args`指定的自定义错误信息。
		135	+
		136	+### SIMD VF编程场景
		137	+
		138	+- `fmt`可以直接使用字符串字面量。使用字符串指针时，该指针必须指向UB中的字符串，示例如下：
		139	+
		140	+ ```cpp
		141	+ __ubuf__ const char* fmt = "value is %d.\n";
		142	+ assert(value == 6, fmt, value);
		143	+ ```
		144	+
		145	+- 每个AIV核在单次`asc_vf_call`执行期间使用2KB预留UB空间临时保存调测数据。同一次`asc_vf_call`中的`assert`、`ascendc_assert`、`printf`和`asc_dump`共享该空间。该空间中的数据传输完成后会被复用，因此上述接口产生的累计调测数据可以超过2KB。单条断言失败信息必须能完整保存在该空间中，否则该条信息不会打印，但接口仍会触发异常。
		146	+- SIMD VF场景下，本接口需要使用2KB预留UB空间传递断言信息。开启`--cce-disable-asc-reserved-ubuf`选项后，本接口不可用。
		147	+- SIMD VF场景下，`simd_printf_fifo_size_per_core`建议配置为3KB以上。配置过小且打印数据量较大时，断言失败信息不会被打印，但接口仍会触发异常。
		148	+
		149	+## 调用示例
		150	+
		151	+### SIMD编程场景
		152	+
		153	+条件不满足时，接口打印AIC/AIV类型、Block号、源码位置、函数名和条件表达式。传入格式字符串和对应参数时，还会打印自定义错误信息。随后，接口触发异常。
		154	+
		155	+以下示例中，`value`为7，不满足`value == 6`，因此触发断言。
87		156
88	```cpp	157	```cpp
89	-#include "utils/debug/asc_assert.h"	158	+extern "C" __global__ __vector__ void AssertSimdKernel()
90	-```
91	-
92	-## SIMD调用示例
93	-
94	-```cpp
95	-// SIMD
96	-__global__ __cube__ void simp_test_equal(int a)
97	{	159	{
98	- // input a is 7	160	+ int32_t value = 7;
99	- assert(a == 6);	161	+ AscendC::printf("before assert, value is %d.\n", value);
		162	+ assert(value == 6, "value is %d.\n", value);
100	}	163	}
101	```	164	```
102		165
103	-程序运行时会触发assert，打印效果如下：	166	+程序运行时会触发断言，实际打印如下：
104		167
105	-```	168	+```plain
106	-[ASSERT] test_assert.asc:41: : Assertion `a == 6' failed.	169	+[AIV Block 0/1] before assert, value is 7.
		170	+[AIV Block 0/1] [ASSERT] .../assert_simd.asc:10: void AssertSimdKernel(): Assertion 'value == 6' failed. value is 7.
107	```	171	```
108		172
109	-## SIMT调用示例	173	+### SIMT VF编程场景
		174	+
		175	+条件不满足时，接口打印源码位置、函数名和条件表达式，然后触发异常。
110		176
111	- SIMT编程场景：	177	- SIMT编程场景：
112		178
	@@ -120,7 +186,7 @@ __global__ __cube__ void simp_test_equal(int a)
120		186
121	程序运行时会触发assert，打印效果如下：	187	程序运行时会触发assert，打印效果如下：
122		188
123	- ```	189	+ ```plain
124	[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(float *): Assertion `!isnan(x[idx])' failed.	190	[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(float *): Assertion `!isnan(x[idx])' failed.
125	```	191	```
126		192
	@@ -136,6 +202,24 @@ __global__ __cube__ void simp_test_equal(int a)
136		202
137	程序运行时会触发assert，打印效果如下：	203	程序运行时会触发assert，打印效果如下：
138		204
139	- ```	205	+ ```plain
140	[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(__gm__ float *): Assertion `!isnan(x[idx])' failed.	206	[ASSERT] /home/.../simt_kernel.asc:44: void simt_kernel(__gm__ float *): Assertion `!isnan(x[idx])' failed.
141	```	207	```
		208	+
		209	+### SIMD VF编程场景
		210	+
		211	+条件不成立时，接口打印源码位置、函数名和条件表达式；如果传入了格式字符串和对应参数，自定义错误信息会追加在该断言失败信息之后。随后，接口触发异常。
		212	+
		213	+```cpp
		214	+__simd_vf__ inline void AssertSimdVf()
		215	+{
		216	+ int32_t value = 7;
		217	+ assert(value == 6, "value is %d.\n", value);
		218	+}
		219	+```
		220	+
		221	+程序运行时会触发断言，打印格式如下：
		222	+
		223	+```plain
		224	+[ASSERT] .../assert_simd_vf.asc:12: void AssertSimdVf(): Assertion 'value == 6' failed. value is 7.
		225	+```

Mdocs/zh/api/Utils-API/tuning_interface/printf.md+2-2

	@@ -77,7 +77,6 @@ static __attribute__((noinline)) void printf(const char* fmt, Args&&... args);
77		77
78	## 约束说明	78	## 约束说明
79		79
80	-- 本接口不支持打印除换行符之外的其他转义字符。
81	<!-- npu="950,A3,910b" id14 -->	80	<!-- npu="950,A3,910b" id14 -->
82	- SIMT场景的支持情况如下：	81	- SIMT场景的支持情况如下：
83	<!-- npu="950" id11 -->	82	<!-- npu="950" id11 -->
	@@ -106,8 +105,9 @@ static __attribute__((noinline)) void printf(const char* fmt, Args&&... args);
106	```	105	```
107		106
108	<!-- npu="950" id15 -->	107	<!-- npu="950" id15 -->
109	-- 在`simd_vf`场景下，每个AIV核使用2KB预留UB空间作为FIFO临时缓冲区。同一次`asc_vf_call`中，`printf`、`asc_dump`和`assert`产生的累计调测数据可以超过2KB，但每条完整编码数据必须能放入该缓冲区；否则，该条数据不会打印。	108	+- SIMD VF场景下，每个AIV核在单次`asc_vf_call`执行期间使用2KB预留UB空间临时保存调测数据。同一次`asc_vf_call`中的`assert`、`ascendc_assert`、`printf`和`asc_dump`共享该空间。该空间中的数据传输完成后会被复用，因此上述接口产生的累计调测数据可以超过2KB。单条调测数据必须能完整保存在该空间中，否则该条数据不会打印。
110	- 每次调用`simd_vf`的`printf`时，除格式字符串和参数外，还会固定占用32字节；整条打印数据需要按8字节对齐。	109	- 每次调用`simd_vf`的`printf`时，除格式字符串和参数外，还会固定占用32字节；整条打印数据需要按8字节对齐。
		110	+- SIMD VF场景下，"simd\_printf\_fifo\_size\_per\_core"建议配置为3KB以上。配置过小且打印数据量较大时，部分调测数据不会被打印。
111	<!-- end id15 -->	111	<!-- end id15 -->
112		112
113	## SIMD调用示例	113	## SIMD调用示例

Mdocs/zh/api/Utils-API/tuning_interface/tuning_interface.md+1-1

	@@ -2,7 +2,7 @@
2		2
3	- [printf](printf.md)	3	- [printf](printf.md)
4		4
5	-- [assert](assert.md)	5	+- [assert/ascendc\_assert](assert.md)
6		6
7	- [\_\_trap](__trap.md)	7	- [\_\_trap](__trap.md)
8		8

Mdocs/zh/api/Utils-API/utils_api_list.md+1-1

	@@ -84,7 +84,7 @@
84	\| 接口名 \| 功能描述 \|	84	\| 接口名 \| 功能描述 \|
85	\| --- \| --- \|	85	\| --- \| --- \|
86	\| [printf](tuning_interface/printf.md) \| 在算子核函数（Kernel）侧的实现代码中，需要输出日志信息时，调用printf接口打印相关内容。 \|	86	\| [printf](tuning_interface/printf.md) \| 在算子核函数（Kernel）侧的实现代码中，需要输出日志信息时，调用printf接口打印相关内容。 \|
87	-\| [assert](tuning_interface/assert.md) \| 本接口在SIMT VF调试场景下提供assert断言功能。在算子核函数（Kernel）侧的SIMT VF实现代码中，如果assert的内部条件判断不为真，则会输出assert条件，并将输入的信息格式化打印在屏幕上。 \|	87	+\| [assert/ascendc_assert](tuning_interface/assert.md) \| 用于在算子核函数（Kernel）中检查指定条件。条件不成立时，接口打印断言失败信息并触发异常。 \|
88	\| [__trap](tuning_interface/__trap.md) \| 在SIMT VF实现代码中调用此接口会中断算子的运行。 \|	88	\| [__trap](tuning_interface/__trap.md) \| 在SIMT VF实现代码中调用此接口会中断算子的运行。 \|
89	\| [clock](tuning_interface/clock.md) \| 本接口在SIMT VF调试场景中提供Clock时间戳功能，用于记录从程序启动到接口调用时刻所经历的时钟周期数（Cycle Count），便于精确分析执行延迟和性能瓶颈。 \|	89	\| [clock](tuning_interface/clock.md) \| 本接口在SIMT VF调试场景中提供Clock时间戳功能，用于记录从程序启动到接口调用时刻所经历的时钟周期数（Cycle Count），便于精确分析执行延迟和性能瓶颈。 \|
90	\| [asc\_dump](tuning_interface/asc_dump.md) \| 将对应内存上的数据打印出来，同时支持打印自定义的附加信息（仅支持uint32\_t类型的信息），比如打印当前行号等。 \|	90	\| [asc\_dump](tuning_interface/asc_dump.md) \| 将对应内存上的数据打印出来，同时支持打印自定义的附加信息（仅支持uint32\_t类型的信息），比如打印当前行号等。 \|

Mdocs/zh/guide/programming_guide/appendix/show_kernel_debug_data_tool.md+1-1

	@@ -1,6 +1,6 @@
1	# show\_kernel\_debug\_data工具	1	# show\_kernel\_debug\_data工具
2		2
3	-在Ascend C算子程序代码中，用户可以使用AscendC::[DumpTensor](../../../api/SIMD-API/basic_api/debug_interface/onboard_print/DumpTensor.md)、AscendC::[printf](../../../api/Utils-API/tuning_interface/printf.md)、AscendC::[PrintTimeStamp](../../../api/SIMD-API/basic_api/debug_interface/onboard_print/PrintTimeStamp.md)、[ascendc\_assert](../../../api/SIMD-API/basic_api/debug_interface/exception_detection/ascendc_assert.md)接口打印相关调试信息，并通过aclInit接口或直接配置acl.json文件，启用Dump配置，导出Ascend C算子核函数（Kernel）的调测信息。本工具提供了对调测信息的离线解析能力，帮助用户获取并解析调试信息，即将导出的bin文件解析成可读格式。本工具的使用示例可参考[show\_kernel\_debug\_data样例](https://gitcode.com/cann/asc-tools/tree/master/examples/01_show_kernel_debug_data)。	3	+在Ascend C算子程序代码中，用户可以使用AscendC::[DumpTensor](../../../api/SIMD-API/basic_api/debug_interface/onboard_print/DumpTensor.md)、AscendC::[printf](../../../api/Utils-API/tuning_interface/printf.md)、AscendC::[PrintTimeStamp](../../../api/SIMD-API/basic_api/debug_interface/onboard_print/PrintTimeStamp.md)、[ascendc\_assert](../../../api/Utils-API/tuning_interface/assert.md)接口打印相关调试信息，并通过aclInit接口或直接配置acl.json文件，启用Dump配置，导出Ascend C算子核函数（Kernel）的调测信息。本工具提供了对调测信息的离线解析能力，帮助用户获取并解析调试信息，即将导出的bin文件解析成可读格式。本工具的使用示例可参考[show\_kernel\_debug\_data样例](https://gitcode.com/cann/asc-tools/tree/master/examples/01_show_kernel_debug_data)。
4		4
5	>[!NOTE]说明	5	>[!NOTE]说明
6	>show\_kernel\_debug\_data支持多用户并发调用，但用户需要指定不同的落盘路径，否则可能出现落盘内容被覆盖等问题。	6	>show\_kernel\_debug\_data支持多用户并发调用，但用户需要指定不同的落盘路径，否则可能出现落盘内容被覆盖等问题。

Mexamples/01_simd_cpp_api/01_utilities/09_log/README.md+1-1

	@@ -134,7 +134,7 @@
134	rtDeviceSynchronize execution failed, reason=aicore exception	134	rtDeviceSynchronize execution failed, reason=aicore exception
135	...	135	...
136		136
137	- [ASSERT] <your_path>/09_log/log.asc:<line>: : Assertion `M < 100 && "M exceeds expected limit!"' failed.	137	+ [ASSERT] <your_path>/09_log/log.asc:<line>: : Assertion 'M < 100 && "M exceeds expected limit!"' failed.
138	[INFO] Execution completed. Check for error messages above.	138	[INFO] Execution completed. Check for error messages above.
139	```	139	```
140		140

Mexamples/01_simd_cpp_api/01_utilities/09_log/README_en.md+1-1

	@@ -134,7 +134,7 @@ Run the following steps in the root directory of this example to build and execu
134	rtDeviceSynchronize execution failed, reason=aicore exception	134	rtDeviceSynchronize execution failed, reason=aicore exception
135	...	135	...
136		136
137	- [ASSERT] <your_path>/09_log/log.asc:<line>: : Assertion `M < 100 && "M exceeds expected limit!"' failed.	137	+ [ASSERT] <your_path>/09_log/log.asc:<line>: : Assertion 'M < 100 && "M exceeds expected limit!"' failed.
138	[INFO] Execution completed. Check for error messages above.	138	[INFO] Execution completed. Check for error messages above.
139	```	139	```
140		140

Mimpl/utils/base/helpers_impl.h+2-1

	@@ -37,7 +37,8 @@ __aicore__ static inline void asc_vf_call_impl(Args&&... args)
37	#ifdef ASCENDC_INTERNAL_SIMD_VF_OVERFLOW_WARNING_ENABLED	37	#ifdef ASCENDC_INTERNAL_SIMD_VF_OVERFLOW_WARNING_ENABLED
38	if (transferFailed) {	38	if (transferFailed) {
39	__asc_aicore::printf_impl(	39	__asc_aicore::printf_impl(
40	- "[WARNING]: SIMD VF debug record exceeds the 2024-byte limit. The record was dropped.\n");	40	+ "[WARNING]: SIMD VF debug output dropped: record exceeds the 2048-byte reserved UB, or "
		41	+ "simd_printf_fifo_size_per_core is too small.\n");
41	}	42	}
42	#else	43	#else
43	(void)transferFailed;	44	(void)transferFailed;

Mimpl/utils/debug/asc_aicore_assert_impl.h+1-1

	@@ -23,7 +23,7 @@ namespace __asc_aicore {
23	inline __aicore__ void __assert_fail(	23	inline __aicore__ void __assert_fail(
24	const __gm__ char* __assertion, const __gm__ char* __file, unsigned int __line, const __gm__ char* __function)	24	const __gm__ char* __assertion, const __gm__ char* __file, unsigned int __line, const __gm__ char* __function)
25	{	25	{
26	- printf_impl_assert("[ASSERT] %s:%u: %s: Assertion `%s' failed.\n", __file, __line, __function, __assertion);	26	+ printf_impl_assert("[ASSERT] %s:%u: %s: Assertion '%s' failed.\n", __file, __line, __function, __assertion);
27	trap();	27	trap();
28	}	28	}
29		29

Mimpl/utils/debug/asc_aicore_printf_impl.h+3-3

	@@ -257,11 +257,11 @@ static __attribute__((noinline)) __aicore__ void printf_impl_assert_msg(
257	set_atomic_none();	257	set_atomic_none();
258	#ifdef __DAV_VEC__	258	#ifdef __DAV_VEC__
259	scalar_printf_impl(	259	scalar_printf_impl(
260	- DumpType::DUMP_ASSERT, fmt, "[AIV Block %u/%u] [ASSERT] %s:%u: %s: Assertion `%s' failed. ",	260	+ DumpType::DUMP_ASSERT, fmt, "[AIV Block %u/%u] [ASSERT] %s:%u: %s: Assertion '%s' failed. ",
261	asc_debug_get_block_idx(), asc_debug_get_block_total_num(), __file, __line, __function, __assertion, args...);	261	asc_debug_get_block_idx(), asc_debug_get_block_total_num(), __file, __line, __function, __assertion, args...);
262	#else	262	#else
263	scalar_printf_impl(	263	scalar_printf_impl(
264	- DumpType::DUMP_ASSERT, fmt, "[AIC Block %u/%u] [ASSERT] %s:%u: %s: Assertion `%s' failed. ",	264	+ DumpType::DUMP_ASSERT, fmt, "[AIC Block %u/%u] [ASSERT] %s:%u: %s: Assertion '%s' failed. ",
265	asc_debug_get_block_idx(), asc_debug_get_block_total_num(), __file, __line, __function, __assertion, args...);	265	asc_debug_get_block_idx(), asc_debug_get_block_total_num(), __file, __line, __function, __assertion, args...);
266	#endif	266	#endif
267	set_ctrl(ctrlValue);	267	set_ctrl(ctrlValue);
	@@ -313,7 +313,7 @@ static __attribute__((noinline)) __aicore__ void printf_impl_assert_msg(
313	const __gm__ char* fmt, Args&&... args)	313	const __gm__ char* fmt, Args&&... args)
314	{	314	{
315	#if !(defined(ASCENDC_DUMP) && ASCENDC_DUMP == 0)	315	#if !(defined(ASCENDC_DUMP) && ASCENDC_DUMP == 0)
316	- std::printf("[ASSERT] %s:%u: %s: Assertion `%s' failed. ", __file, __line, __function, __assertion);	316	+ std::printf("[ASSERT] %s:%u: %s: Assertion '%s' failed. ", __file, __line, __function, __assertion);
317	std::printf(fmt, args...);	317	std::printf(fmt, args...);
318	#endif	318	#endif
319	}	319	}

Mimpl/utils/debug/asc_assert_simd_impl.h+2-1

	@@ -41,12 +41,13 @@ static __attribute__((noinline)) __simd_callee__ void __assert_fail(
41	__ubuf__ const char* assertion, __ubuf__ const char* file, unsigned int line,	41	__ubuf__ const char* assertion, __ubuf__ const char* file, unsigned int line,
42	__ubuf__ const char* function) noexcept	42	__ubuf__ const char* function) noexcept
43	{	43	{
44	- printf_impl_assert("[ASSERT] %s:%u: %s: Assertion `%s' failed.\n", file, line, function, assertion);	44	+ printf_impl_assert("[ASSERT] %s:%u: %s: Assertion '%s' failed.\n", file, line, function, assertion);
45	wait_vf_assert_handshake();	45	wait_vf_assert_handshake();
46	__trap();	46	__trap();
47	}	47	}
48		48
49	template <typename... Args>	49	template <typename... Args>
		50	+__ASC_USE_RESERVED_UBUF__(3510, "assert is forbidden when compile option --cce-disable-asc-reserved-ubuf is enabled")
50	static __attribute__((noinline)) __simd_callee__ void __assert_fail_msg(	51	static __attribute__((noinline)) __simd_callee__ void __assert_fail_msg(
51	__ubuf__ const char* assertion, __ubuf__ const char* file, unsigned int line, __ubuf__ const char* function,	52	__ubuf__ const char* assertion, __ubuf__ const char* file, unsigned int line, __ubuf__ const char* function,
52	__ubuf__ const char* fmt, Args&&... args) noexcept	53	__ubuf__ const char* fmt, Args&&... args) noexcept

Mimpl/utils/debug/asc_debug_utils.h+4-1

	@@ -205,6 +205,9 @@ __aicore__ inline void do_overlow_skip(
205	__aicore__ inline bool ringbuf_overflow_wait(__gm__ DebugBlockReadInfo* readInfo, __gm__ DebugBlockWriteInfo* writeInfo)	205	__aicore__ inline bool ringbuf_overflow_wait(__gm__ DebugBlockReadInfo* readInfo, __gm__ DebugBlockWriteInfo* writeInfo)
206	{	206	{
207	constexpr uint32_t maxCounter = 15;	207	constexpr uint32_t maxCounter = 15;
		208	+ // DAV_3510 SYS_CNT runs at 1 GHz, 20 times the 50 MHz baseline used by the default timeout.
		209	+ // Scale the cycle budget to retain a 300 ms wait and give RTS time to consume pending ring-buffer data.
		210	+ constexpr uint64_t rtsWaitCycle = 15 * 1000 * 1000 * 20;
208	uint32_t counter = 0;	211	uint32_t counter = 0;
209	volatile uint64_t readofst = ld_dev((__gm__ uint64_t*)(&(readInfo->bufOffset)), 0);	212	volatile uint64_t readofst = ld_dev((__gm__ uint64_t*)(&(readInfo->bufOffset)), 0);
210	volatile uint64_t writeofst = ld_dev((__gm__ uint64_t*)(&(writeInfo->bufOffset)), 0);	213	volatile uint64_t writeofst = ld_dev((__gm__ uint64_t*)(&(writeInfo->bufOffset)), 0);
	@@ -213,7 +216,7 @@ __aicore__ inline bool ringbuf_overflow_wait(__gm__ DebugBlockReadInfo* readInfo
213	if (counter >= maxCounter) { // max wait 15 * 300ms, rts read gm per 200ms	216	if (counter >= maxCounter) { // max wait 15 * 300ms, rts read gm per 200ms
214	return false;	217	return false;
215	}	218	}
216	- ringbuf_wait_rts_sync(); // wait 20 * 15 ms	219	+ ringbuf_wait_rts_sync<rtsWaitCycle>();
217	++counter;	220	++counter;
218	asc_entire_dcci((__gm__ uint64_t*)readInfo);	221	asc_entire_dcci((__gm__ uint64_t*)readInfo);
219	readofst = ld_dev((__gm__ uint64_t*)(&(readInfo->bufOffset)), 0);	222	readofst = ld_dev((__gm__ uint64_t*)(&(readInfo->bufOffset)), 0);

Mimpl/utils/debug/asc_simd_printf_impl.h+1-1

	@@ -248,7 +248,7 @@ __simd_callee__ inline void printf_impl_assert_msg(
248	#if !(defined(ASCENDC_DUMP) && ASCENDC_DUMP == 0)	248	#if !(defined(ASCENDC_DUMP) && ASCENDC_DUMP == 0)
249	enable_asc_assert();	249	enable_asc_assert();
250	scalar_printf_impl(	250	scalar_printf_impl(
251	- DumpType::DUMP_ASSERT, fmt, "[ASSERT] %s:%u: %s: Assertion `%s' failed. ", file, line, function, assertion,	251	+ DumpType::DUMP_ASSERT, fmt, "[ASSERT] %s:%u: %s: Assertion '%s' failed. ", file, line, function, assertion,
252	args...);	252	args...);
253	#endif	253	#endif
254	}	254	}

Mimpl/utils/debug/npu_arch_3510/asc_debug_utils_impl.h+34-9

	@@ -49,7 +49,8 @@ __simd_callee__ constexpr inline uint32_t align_up(uint32_t a, uint32_t b) { ret
49		49
50	__simd_callee__ inline void wait_vf_debug_buffer_drained(__ubuf__ BlockVFBufInfo* block_info)	50	__simd_callee__ inline void wait_vf_debug_buffer_drained(__ubuf__ BlockVFBufInfo* block_info)
51	{	51	{
52	- while (block_info->readLen != block_info->writeLen) {	52	+ // flag != 0: the consumer stopped draining, readLen will never catch up.
		53	+ while (block_info->readLen != block_info->writeLen && block_info->flag == 0) {
53	__asm__ __volatile__("");	54	__asm__ __volatile__("");
54	}	55	}
55	}	56	}
	@@ -127,9 +128,16 @@ __simd_callee__ inline void wait_vf_assert_handshake()
127	{	128	{
128	__ubuf__ BlockVFBufInfo* block_info = get_printf_ubuf_addr(0);	129	__ubuf__ BlockVFBufInfo* block_info = get_printf_ubuf_addr(0);
129	wait_vf_debug_buffer_drained(block_info);	130	wait_vf_debug_buffer_drained(block_info);
		131	+ // Nobody left to acknowledge the handshake; drop the message but still let the caller trap.
		132	+ if (block_info->flag != 0) {
		133	+ return;
		134	+ }
130	__ubuf__ volatile BlockVFBufInfo::AssertState* assertFlag = &block_info->assertFlag;	135	__ubuf__ volatile BlockVFBufInfo::AssertState* assertFlag = &block_info->assertFlag;
131	*assertFlag = BlockVFBufInfo::AssertState::RAISED;	136	*assertFlag = BlockVFBufInfo::AssertState::RAISED;
132	while (*assertFlag != BlockVFBufInfo::AssertState::DRAINED) {	137	while (*assertFlag != BlockVFBufInfo::AssertState::DRAINED) {
		138	+ if (block_info->flag != 0) {
		139	+ break;
		140	+ }
133	__asm__ __volatile__("");	141	__asm__ __volatile__("");
134	}	142	}
135	}	143	}
	@@ -321,6 +329,14 @@ __aicore__ inline void asc_vf_debug_publish(
321		329
322	__gm__ BlockRingBufInfo* blockRingBufInfo = get_block_ring_buf_info();	330	__gm__ BlockRingBufInfo* blockRingBufInfo = get_block_ring_buf_info();
323	auto* debugBlockInfo = reinterpret_cast<__gm__ DebugBlockHeadInfo*>(blockRingBufInfo);	331	auto* debugBlockInfo = reinterpret_cast<__gm__ DebugBlockHeadInfo*>(blockRingBufInfo);
		332	+ // A batch larger than the ring buffer can never be published; retrying would spin forever and hang the kernel.
		333	+ // Drop it: advance readLen to release the producer, raise flag to stop the transfer loop.
		334	+ if (tlvLen > blockRingBufInfo->ringBufLen \|\| sizeof(SkipTlv) >= blockRingBufInfo->ringBufLen) {
		335	+ blockInfo->readLen = curWriteLen;
		336	+ blockInfo->flag = 1;
		337	+ return;
		338	+ }
		339	+ // Remaining failures are transient (RTS lagging); keep retrying.
324	if (!check_ringbuf_space(debugBlockInfo, tlvLen)) {	340	if (!check_ringbuf_space(debugBlockInfo, tlvLen)) {
325	return;	341	return;
326	}	342	}
	@@ -346,13 +362,16 @@ __aicore__ inline void asc_vf_debug_publish(
346	__aicore__ inline bool asc_vf_debug_ub2gm()	362	__aicore__ inline bool asc_vf_debug_ub2gm()
347	{	363	{
348	__ubuf__ BlockVFBufInfo* blockInfo = get_printf_ubuf_addr_aicore(0);	364	__ubuf__ BlockVFBufInfo* blockInfo = get_printf_ubuf_addr_aicore(0);
		365	+ // The VF unit updates these concurrently; without volatile the loads are hoisted and the loop
		366	+ // works off a stale snapshot (finish == 1 while writeLen still reads 0), dropping every record.
		367	+ __ubuf__ volatile BlockVFBufInfo* vInfo = blockInfo;
349	for (;;) {	368	for (;;) {
350	- uint32_t curReadLen = blockInfo->readLen;	369	+ uint32_t curReadLen = vInfo->readLen;
351	- uint32_t curWriteLen = blockInfo->writeLen;	370	+ uint32_t curWriteLen = vInfo->writeLen;
352		371
353	- const bool isValidHeader = blockInfo->magic == ASCENDC_SIMD_VF_MAGIC_NUMBER &&	372	+ const bool isValidHeader = vInfo->magic == ASCENDC_SIMD_VF_MAGIC_NUMBER &&
354	- blockInfo->length <= ASCENDC_SIMD_VF_PRINTF_UBUF_MAX_SIZE &&	373	+ vInfo->length <= ASCENDC_SIMD_VF_PRINTF_UBUF_MAX_SIZE &&
355	- curWriteLen <= blockInfo->length;	374	+ curWriteLen <= vInfo->length;
356	if (!isValidHeader) {	375	if (!isValidHeader) {
357	blockInfo->flag = 1;	376	blockInfo->flag = 1;
358	break;	377	break;
	@@ -370,12 +389,18 @@ __aicore__ inline bool asc_vf_debug_ub2gm()
370	continue;	389	continue;
371	}	390	}
372		391
373	- if (blockInfo->assertFlag == BlockVFBufInfo::AssertState::RAISED) {	392	+ if (vInfo->assertFlag == BlockVFBufInfo::AssertState::RAISED) {
374	- blockInfo->assertFlag = BlockVFBufInfo::AssertState::DRAINED;	393	+ vInfo->assertFlag = BlockVFBufInfo::AssertState::DRAINED;
375	break;	394	break;
376	}	395	}
377		396
378	- if (blockInfo->finish == 1) {	397	+ if (vInfo->finish == 1) {
		398	+ // finish is set on the scalar unit; the VF writes may not be visible yet. Sync and re-check
		399	+ // before concluding the buffer is empty, otherwise a pending record is silently dropped.
		400	+ pipe_barrier(PIPE_ALL);
		401	+ if (vInfo->readLen < vInfo->writeLen) {
		402	+ continue;
		403	+ }
379	break;	404	break;
380	}	405	}
381	}	406	}

	@@ -1,9 +1,5 @@
1	# 异常检测<a name="ZH-CN_TOPIC_0000002512364579"></a>	1	# 异常检测<a name="ZH-CN_TOPIC_0000002512364579"></a>
2		2
3	-- [ascendc\_assert](ascendc_assert.md)
4	-
5	-- [assert](assert.md)
6	-
7	- [Trap](Trap.md)	3	- [Trap](Trap.md)
8		4
9	- [CheckLocalMemoryIA\(ISASI\)](CheckLocalMemoryIA_ISASI.md)	5	- [CheckLocalMemoryIA\(ISASI\)](CheckLocalMemoryIA_ISASI.md)

	@@ -2,7 +2,7 @@
2		2
3	- [printf](printf.md)	3	- [printf](printf.md)
4		4
5	-- [assert](assert.md)	5	+- [assert/ascendc\_assert](assert.md)
6		6
7	- [\_\_trap](__trap.md)	7	- [\_\_trap](__trap.md)
8		8