SIMT 增加assert打印接口-asc-devkit-AtomGit

已合并

SIMT 增加assert打印接口 #1565

ruoshuisixue创建于 4月13日从ruoshuisixue/asc-devkit_doc： cherry-pick-mr-1564-1776053255186-auto 合入到 cann/asc-devkit： master

SIMT 增加assert打印接口 #1565

已合并

ruoshuisixue创建于 4月13日

从ruoshuisixue/asc-devkit_doc cherry-pick-mr-1564-1776053255186-auto合入到cann/asc-devkit master

共 4 个文件变更+204-4

Mexamples/03_simt_api/01_utilities/00_printf/printf.asc+2-2

	@@ -10,7 +10,7 @@
10		10
11		11
12	/* !	12	/* !
13	- * \file gather.asc	13	+ * \file printf.asc
14	* \brief	14	* \brief
15	*/	15	*/
16		16
	@@ -20,7 +20,7 @@
20	#include "asc_simt.h"	20	#include "asc_simt.h"
21	#include "asc_printf.h"	21	#include "asc_printf.h"
22		22
23	-__global__ void simt_printf(float* input, uint32_t in_shape)	23	+__global__ void simt_printf(float* input, uint32_t in_shape)
24	{	24	{
25	// Calculate global thread ID	25	// Calculate global thread ID
26	int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;	26	int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;

Aexamples/03_simt_api/01_utilities/01_assert/CMakeLists.txt+29-0

@@ -0,0 +1,29 @@
	1	+# ----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------------------------------------
	10	+
	11	+cmake_minimum_required(VERSION 3.16)
	12	+
	13	+find_package(ASC REQUIRED)
	14	+
	15	+project(kernel_samples LANGUAGES ASC CXX)
	16	+
	17	+add_executable(demo
	18	+ assert.asc
	19	+)
	20	+
	21	+# ======================================================================================
	22	+# NPU 编译选项配置
	23	+#
	24	+# 说明：
	25	+# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
	26	+# ======================================================================================
	27	+target_compile_options(demo PRIVATE
	28	+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510 --enable-simt>
	29	+)

Mexamples/03_simt_api/01_utilities/01_assert/README.md+91-2

	@@ -1,3 +1,92 @@
1	-# assert样例介绍	1	+# 基于gather算子的SIMT assert断言功能实现样例
2		2
3	-待开发。	3	+## 概述
		4	+
		5	+本样例演示在SIMT编程下使用```assert()```接口实现上板进行功能调试的方法。
		6	+
		7	+## 支持的产品
		8	+
		9	+- Ascend 950PR/Ascend 950DT
		10	+
		11	+## 目录结构
		12	+
		13	+```
		14	+├── 01_assert
		15	+│ ├── CMakeLists.txt # cmake编译文件
		16	+│ ├── assert.asc # Ascend C算子实现加assert断言的调用样例
		17	+\| └── README.md
		18	+```
		19	+
		20	+## 算子描述
		21	+
		22	+- 算子功能:
		23	+
		24	+ 本样例详细展示了在SIMT实现函数中使用```assert()```接口的实践方式，实现对算子执行过程中断言的调试。
		25	+
		26	+
		27	+- 算子实现:
		28	+ ```cpp
		29	+ __global__ void simt_assert(float* input, uint32_t in_shape)
		30	+ {
		31	+ // Calculate global thread ID
		32	+ int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
		33	+ if (threadIdx.x < 1) {
		34	+ printf("[SIMT] %s\n", "trap check start 1!");
		35	+ printf("[SIMT] %s\n", "trap check start 2!");
		36	+ printf("[SIMT] %s\n", "trap check start 3!");
		37	+ assert(in_shape < 1);
		38	+ printf("[SIMT] %s\n", "trap check 1!");
		39	+ } else if(threadIdx.x < 5) {
		40	+ printf("[SIMT] %s\n", "trap check 2!");
		41	+ assert(in_shape > 1);
		42	+ printf("[SIMT] %s\n", "trap check 3!");
		43	+ }
		44	+ }
		45	+ ```
		46	+
		47	+## 编译运行
		48	+
		49	+在本样例根目录下执行如下步骤，编译并执行算子。
		50	+
		51	+- 配置环境变量
		52	+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
		53	+ - 默认路径，root用户安装CANN软件包
		54	+ ```bash
		55	+ source /usr/local/Ascend/cann/set_env.sh
		56	+ ```
		57	+
		58	+ - 默认路径，非root用户安装CANN软件包
		59	+ ```bash
		60	+ source $HOME/Ascend/cann/set_env.sh
		61	+ ```
		62	+
		63	+ - 指定路径install_path，安装CANN软件包
		64	+ ```bash
		65	+ source ${install_path}/cann/set_env.sh
		66	+ ```
		67	+
		68	+- 样例执行
		69	+ ```bash
		70	+ mkdir -p build && cd build; # 创建并进入build目录
		71	+ cmake ..; make -j; # 编译工程
		72	+ ./demo # 执行样例
		73	+ ```
		74	+ 执行后有如下打印信息，说明功能正常。
		75	+ ```
		76	+ [SIMT] trap check 2!
		77	+ [SIMT] trap check 2!
		78	+ [SIMT] trap check 2!
		79	+ [SIMT] trap check 2!
		80	+ [SIMT] trap check 2!
		81	+ [SIMT] trap check 2!
		82	+ [SIMT] trap check 2!
		83	+ [SIMT] trap check 2!
		84	+ [SIMT] trap check start 1!
		85	+ [SIMT] trap check start 1!
		86	+ [SIMT] trap check start 2!
		87	+ [SIMT] trap check start 2!
		88	+ [SIMT] trap check start 3!
		89	+ [SIMT] trap check start 3!
		90	+ [ASSERT] xxx/assert.asc:32: void simt_assert(float *, uint32_t): Assertion `in_shape < 1' failed.
		91	+ [ASSERT] xxx/assert.asc:32: void simt_assert(float *, uint32_t): Assertion `in_shape < 1' failed.
		92	+ ```

Aexamples/03_simt_api/01_utilities/01_assert/assert.asc+82-0

@@ -0,0 +1,82 @@
	1	+/**
	2	+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+* CANN Open Software License Agreement Version 2.0 (the "License").
	5	+* Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+* See LICENSE in the root of the software repository for the full text of the License.
	9	+*/
	10	+
	11	+
	12	+/* !
	13	+ * \file assert.asc
	14	+ * \brief
	15	+ */
	16	+
	17	+#include <iostream>
	18	+#include <vector>
	19	+#include "acl/acl.h"
	20	+#include "asc_simt.h"
	21	+#include "asc_printf.h"
	22	+#include "asc_assert.h"
	23	+
	24	+__global__ void simt_assert(float* input, uint32_t in_shape)
	25	+{
	26	+ // Calculate global thread ID
	27	+ int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
	28	+ if (threadIdx.x < 1) {
	29	+ printf("[SIMT] %s\n", "trap check start 1!");
	30	+ printf("[SIMT] %s\n", "trap check start 2!");
	31	+ printf("[SIMT] %s\n", "trap check start 3!");
	32	+ assert(in_shape < 1);
	33	+ printf("[SIMT] %s\n", "trap check 1!");
	34	+ } else if (threadIdx.x < 5) {
	35	+ printf("[SIMT] %s\n", "trap check 2!");
	36	+ assert(in_shape > 1);
	37	+ printf("[SIMT] %s\n", "trap check 3!");
	38	+ }
	39	+}
	40	+
	41	+void host_func(std::vector<float>& input, const uint32_t in_shape)
	42	+{
	43	+ uint32_t input_total_length = input.size();
	44	+ size_t input_total_byte_size = in_shape * sizeof(float);
	45	+
	46	+ int32_t device_id = 0;
	47	+ aclrtStream stream = nullptr;
	48	+
	49	+ uint8_t* input_host = reinterpret_cast<uint8_t *>(input.data());
	50	+ float* input_device = nullptr;
	51	+ aclInit(nullptr);
	52	+ aclrtSetDevice(device_id);
	53	+ aclrtCreateStream(&stream);
	54	+ aclrtMalloc((void **)&input_device, input_total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
	55	+ aclrtMemcpy(input_device, input_total_byte_size, input_host, input_total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
	56	+
	57	+ uint32_t num_blocks = 2;
	58	+ uint32_t num_threads = 32;
	59	+ uint32_t dyn_ubuf_size = 0; // No need to alloc dynamic memory.
	60	+ simt_assert<<<num_blocks, num_threads, dyn_ubuf_size, stream>>>(input_device, in_shape);
	61	+ aclrtSynchronizeStream(stream);
	62	+
	63	+ aclrtFree(input_device);
	64	+ aclrtDestroyStream(stream);
	65	+ aclrtResetDevice(device_id);
	66	+ aclFinalize();
	67	+
	68	+ return;
	69	+}
	70	+
	71	+
	72	+int32_t main(int32_t argc, char* argv[])
	73	+{
	74	+ constexpr uint32_t in_shape = 128;
	75	+ std::vector<float> input(in_shape);
	76	+ for (uint32_t i = 0; i < in_shape; i++) {
	77	+ input[i] = 1.118f + i;
	78	+ }
	79	+
	80	+ host_func(input, in_shape);
	81	+ return 0;
	82	+}