【PR】:add-common-faq-runtime-AtomGit

已合并

【PR】:add-common-faq #2806

rui创建于 6月13日从good_luck_to_me/runtime： doc/faq-clean 合入到 cann/runtime： master

【PR】:add-common-faq #2806

已合并

rui创建于 6月13日

从good_luck_to_me/runtime doc/faq-clean合入到cann/runtime master

收起

docs
- 04_FAQ
  - ACLGraph捕获过程中任务提交限制.md
  - AI应用进程未退出导致休眠唤醒失败.md
  - Runtime版本与CANN版本不匹配导致的问题.md
  - Stream同步与Event同步的区别与选择.md
  - aclInit初始化失败常见原因排查.md
  - aclrtMalloc内存申请失败常见原因.md
  - aclrtMemcpyAsync在错误的Stream上下发失败.md
  - aclrtSetDevice调用失败.md
  - 低版本内核使用asan导致算子执行失败.md
  - 多Device场景下Stream跨Device下发失败.md
  - 如何理解默认Device和默认Stream机制.md
  - 如何获取和解读Runtime异步错误码.md
  - 如何选择合适的内存分配策略.md
  - 如何通过plog日志定位Device侧异常.md
  - 析构函数中调用aclFinalize导致应用进程coredump.md
  - 用户进程异常退出后重启进程失败.md
  - 算子执行输出全0的常见原因排查.md
  - 跨DeviceP2P数据交互配置失败.md
  - 进程间IPC内存共享的页表对齐要求.md
  - 遇错即停模式下错误定位方法.md
- README.md

共 21 个文件变更+2175-403

Adocs/04_FAQ/ACLGraph捕获过程中任务提交限制.md+114-0

@@ -0,0 +1,114 @@
	1	+# ACL Graph捕获过程中任务提交限制
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：对Stream、Event、Device、Context进行同步或查询操作时失败
	6	+
	7	+使用ACL Graph方式捕获Stream上的任务时（在`aclmdlRICaptureBegin`和`aclmdlRICaptureEnd`之间），对Stream、Event、Device、Context进行同步或查询操作，会导致捕获失败。
	8	+
	9	+错误代码示例：
	10	+```cpp
	11	+aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
	12	+aclrtSynchronizeStream(stream); // ✗ 非法操作，会导致捕获失败
	13	+aclmdlRICaptureEnd(stream, &modelRI);
	14	+```
	15	+
	16	+报错日志示例如下：
	17	+```
	18	+[ERROR] RUNTIME: operation not permitted when a stream is capturing and the specified capture mode is not relaxed, ret=107041
	19	+```
	20	+
	21	+现象2：全局禁止模式下调用内存同步操作类函数时报错
	22	+
	23	+在全局禁止模式（`ACL_MODEL_RI_CAPTURE_MODE_GLOBAL`）下调用内存同步操作类函数（如`aclrtMemset`、`aclrtMemcpy`、`aclrtMemcpy2d`）时报错。
	24	+
	25	+错误代码示例：
	26	+```cpp
	27	+aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
	28	+aclrtMemcpy(dst, size, src, size, ACL_MEMCPY_DEVICE_TO_DEVICE); // ✗ 非法操作，会报错
	29	+aclmdlRICaptureEnd(stream, &modelRI);
	30	+```
	31	+
	32	+报错日志示例如下：
	33	+```
	34	+[ERROR] RUNTIME: operation not permitted when a stream is capturing and the specified capture mode is not relaxed, ret=107041
	35	+```
	36	+
	37	+现象3：使用默认Stream进行捕获导致失败
	38	+
	39	+使用默认Stream（传nullptr）进行ACL Graph捕获时，可能导致捕获失败或行为异常。
	40	+
	41	+错误代码示例：
	42	+```cpp
	43	+aclmdlRICaptureBegin(nullptr, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL); // ✗ 使用默认Stream
	44	+aclmdlRICaptureEnd(nullptr, &modelRI);
	45	+```
	46	+
	47	+## 可能原因
	48	+
	49	+1. 对Stream、Event、Device、Context进行同步或查询操作：捕获阶段内执行同步或查询操作违反了捕获规则。
	50	+2. 全局禁止模式下调用内存同步操作类函数：ACL_MODEL_RI_CAPTURE_MODE_GLOBAL模式下不允许调用aclrtMemset、aclrtMemcpy等同步内存函数。
	51	+3. 使用默认Stream进行捕获：默认Stream（nullptr）在捕获过程中可能导致失败或行为异常。
	52	+
	53	+## 处理步骤
	54	+
	55	+### 原因1：对Stream、Event、Device、Context进行同步或查询操作
	56	+
	57	+在捕获阶段前或捕获阶段后进行同步/查询操作：
	58	+
	59	+```cpp
	60	+// 正确示例：在捕获前完成同步操作
	61	+aclrtSynchronizeStream(stream); // ✓ 捕获前同步
	62	+
	63	+aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
	64	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_DEVICE_TO_DEVICE, stream);
	65	+aclmdlRICaptureEnd(stream, &modelRI);
	66	+
	67	+// 执行模型后进行同步
	68	+aclmdlRIExecuteAsync(modelRI, stream);
	69	+aclrtSynchronizeStream(stream); // ✓ 捕获后同步
	70	+```
	71	+
	72	+### 原因2：全局禁止模式下调用内存同步操作类函数
	73	+
	74	+如果业务确定内存同步操作不会影响任务捕获，可以切换为`ACL_MODEL_RI_CAPTURE_MODE_RELAXED`模式：
	75	+
	76	+```cpp
	77	+aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
	78	+
	79	+// 异步内存复制（始终允许）
	80	+aclrtMemcpyAsync(self_d, size, self_h, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	81	+
	82	+// 切换捕获模式为RELAXED，允许调用aclrtMemcpy函数
	83	+aclmdlRICaptureMode mode = ACL_MODEL_RI_CAPTURE_MODE_RELAXED;
	84	+aclmdlRICaptureThreadExchangeMode(&mode);
	85	+
	86	+// 同步内存复制（仅执行一次，RELAXED模式下允许）
	87	+aclrtMemcpy(other_d, size, other_h, size, ACL_MEMCPY_HOST_TO_DEVICE);
	88	+
	89	+// 将捕获模式切换回GLOBAL
	90	+aclmdlRICaptureThreadExchangeMode(&mode);
	91	+
	92	+aclmdlRICaptureEnd(stream, &modelRI);
	93	+```
	94	+
	95	+### 原因3：使用默认Stream进行捕获
	96	+
	97	+创建并使用显式Stream，而不是使用默认Stream：
	98	+
	99	+```cpp
	100	+// 错误示例：在捕获过程中使用默认Stream
	101	+aclmdlRICaptureBegin(nullptr, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL); // ✗ 使用默认Stream
	102	+aclmdlRICaptureEnd(nullptr, &modelRI);
	103	+
	104	+// 正确示例：创建显式Stream
	105	+aclrtStream stream;
	106	+aclrtCreateStream(&stream);
	107	+aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL); // ✓ 使用显式Stream
	108	+aclmdlRICaptureEnd(stream, &modelRI);
	109	+aclrtDestroyStream(stream);
	110	+```
	111	+
	112	+## 相关 issue
	113	+
	114	+- [Issue #487: ACL Graph捕获阶段控核接口支持性咨询](https://gitcode.com/cann/runtime/issues/487)

Ddocs/04_FAQ/AI应用进程未退出导致休眠唤醒失败.md+0-20

	@@ -1,20 +0,0 @@
1	-# AI应用进程未退出，导致休眠唤醒失败
2	-
3	-## 问题现象描述
4	-
5	-休眠失败。
6	-
7	-查看应用类日志，系统内部的任务分发模块hwts正处于busy状态，检查发现不满足休眠条件，日志片段示例如下：
8	-
9	-```
10	-[ERROR] TSCH(-1,null):2023-01-01-02:53:45.850.781 1 (dieid:0,cpuid:0) device_management_plat.c:563 suspend_ack: suspend pre check fail, hwts is busy
11	-[EVENT] TSCH(-1,null):2023-01-01-02:53:45.850.803 2 (dieid:0,cpuid:0) device_management.c:411 process_low_power_cmd: ts suspend ack ret=1.
12	-```
13	-
14	-## 原因分析
15	-
16	-根据休眠唤醒的流程，休眠前AI应用进程必须先退出，相关硬件资源处于idle态，才允许休眠。不满足休眠条件，会有相关报错，本案例中因为AI应用进程未退出，在休眠唤醒时检测到hwts处于busy状态，因此休眠失败。
17	-
18	-## 解决办法
19	-
20	-用户需要确保AI应用进程已经运行结束或者优雅退出，推荐使用kill -2 _PID_退出相关进程，_PID_需替换为实际进程ID。

Adocs/04_FAQ/Runtime版本与CANN版本不匹配导致的问题.md+93-0

@@ -0,0 +1,93 @@
	1	+# Runtime版本与CANN版本不匹配导致的问题
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：接口返回功能不支持错误码
	6	+
	7	+调用 Runtime 接口时返回 ACL_ERROR_RT_FEATURE_NOT_SUPPORT（错误码 207000），表示当前版本不支持该功能。
	8	+
	9	+报错日志示例如下：
	10	+```
	11	+aclrtGetVersion failed, ret = 207000, feature not supported
	12	+```
	13	+
	14	+现象2：接口行为异常
	15	+
	16	+接口参数不兼容或返回值不符合预期，运行时库版本与编译时版本不一致导致行为差异。
	17	+
	18	+典型场景示例：
	19	+```c
	20	+// 编译时使用新版本头文件，运行时加载旧版本库
	21	+aclError ret = aclrtGetVersion(&major, &minor, &patch); // 返回值与预期不符
	22	+```
	23	+
	24	+现象3：编译链接错误
	25	+
	26	+编译时找不到某些接口定义，头文件版本与库文件版本不匹配。
	27	+
	28	+报错日志示例如下：
	29	+```
	30	+undefined reference to `aclrtGetVersion'
	31	+```
	32	+
	33	+## 可能原因
	34	+
	35	+1. 不同 CANN 版本的 API 差异：新版本引入新接口、旧版本不支持某些功能。
	36	+2. 运行时库版本不一致：编译时链接的库版本与运行时加载的库版本不同。
	37	+3. 环境变量配置错误：ASCEND_HOME、LD_LIBRARY_PATH 等指向错误的版本目录。
	38	+
	39	+## 处理步骤
	40	+
	41	+### 原因1：不同 CANN 版本的 API 差异
	42	+
	43	+解决方法：
	44	+- 检查版本信息：使用 aclrtGetVersion 查询当前 Runtime 版本
	45	+- 参考 API 文档：确认接口在不同版本的支持情况
	46	+- 升级或降级版本：根据需求调整 CANN 版本
	47	+
	48	+版本查询示例：
	49	+```c
	50	+size_t majorVersion = 0;
	51	+size_t minorVersion = 0;
	52	+size_t patchVersion = 0;
	53	+aclError ret = aclrtGetVersion(&majorVersion, &minorVersion, &patchVersion);
	54	+printf("Runtime version: %zu.%zu.%zu\n", majorVersion, minorVersion, patchVersion);
	55	+```
	56	+
	57	+### 原因2：运行时库版本不一致
	58	+
	59	+解决方法：
	60	+- 检查编译链接库：确认编译时链接的 libascendcl.so 版本
	61	+- 检查运行时加载库：使用 `ldd` 命令查看实际加载的库路径
	62	+- 确保版本一致：编译和运行时使用相同的 CANN 版本
	63	+
	64	+命令示例：
	65	+```bash
	66	+# 查看可执行文件链接的库
	67	+ldd your_program \| grep ascendcl
	68	+
	69	+# 查看 CANN 安装版本
	70	+cat /usr/local/Ascend/version.info
	71	+```
	72	+
	73	+### 原因3：环境变量配置错误
	74	+
	75	+解决方法：
	76	+- 检查 ASCEND_HOME：确认指向正确的 CANN 安装目录
	77	+- 检查 LD_LIBRARY_PATH：确保包含正确的库路径
	78	+- 更新环境变量：修改 ~/.bashrc 或 ~/.bash_profile
	79	+
	80	+命令示例：
	81	+```bash
	82	+# 查看环境变量
	83	+echo $ASCEND_HOME
	84	+echo $LD_LIBRARY_PATH
	85	+
	86	+# 设置环境变量（示例）
	87	+export ASCEND_HOME=/usr/local/Ascend
	88	+export LD_LIBRARY_PATH=$ASCEND_HOME/lib64:$LD_LIBRARY_PATH
	89	+```
	90	+
	91	+## 相关 issue
	92	+
	93	+暂无相关Issue。

Adocs/04_FAQ/Stream同步与Event同步的区别与选择.md+139-0

@@ -0,0 +1,139 @@
	1	+# Stream同步与Event同步的区别与选择
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：不理解Stream同步和Event同步的差异
	6	+
	7	+混淆两种同步机制的使用范围和特性，导致选择不当。
	8	+
	9	+典型错误场景：
	10	+```c
	11	+// 需要等待跨Stream任务完成，但使用了Stream同步
	12	+aclrtSynchronizeStream(stream1); // 仅等待 stream1，不等待 stream2
	13	+```
	14	+
	15	+现象2：同步时机选择不当导致性能问题或死锁
	16	+
	17	+错误选择同步方式：单流等待用Event（过度设计），多流协调用Stream同步（会阻塞）。
	18	+
	19	+典型错误场景：
	20	+```c
	21	+// 单Stream内等待，但使用了Event同步（过度设计）
	22	+aclrtEvent event;
	23	+aclrtRecordEvent(event, stream);
	24	+aclrtSynchronizeEvent(event); // Event同步比Stream同步更复杂
	25	+```
	26	+
	27	+现象3：错误使用同步接口导致任务执行顺序混乱
	28	+
	29	+跨Stream协调时未使用Event同步，导致任务依赖关系不正确。
	30	+
	31	+## 可能原因
	32	+
	33	+1. 不理解两种同步机制差异：Stream 同步阻塞整流，Event 同步精细控制。
	34	+2. 错误选择同步方式：单流等待用 Event，多流协调用 Stream，顺序颠倒。
	35	+
	36	+## 处理步骤
	37	+
	38	+### 原因1：不理解两种同步机制差异
	39	+
	40	+解决方法：
	41	+- 理解 Stream 同步特性：
	42	+ - 作用范围：阻塞指定 Stream 上的所有任务，直到全部完成
	43	+ - 适用场景：等待单个 Stream 的整批任务完成
	44	+ - 接口：aclrtSynchronizeStream(stream)
	45	+- 理解 Event 同步特性：
	46	+ - 作用范围：精确控制到某个时间点，支持跨 Stream 同步
	47	+ - 适用场景：跨 Stream 任务协调、精细时间点控制
	48	+ - 接口：aclrtRecordEvent、aclrtSynchronizeEvent、aclrtStreamWaitEvent
	49	+
	50	+对比示例：
	51	+```c
	52	+// Stream 同步：等待整流任务完成
	53	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	54	+myKernel<<<..., stream>>>();
	55	+aclrtMemcpyAsync(host, size, dst, size, ACL_MEMCPY_DEVICE_TO_HOST, stream);
	56	+
	57	+aclrtSynchronizeStream(stream); // 等待上述3个任务全部完成
	58	+
	59	+// Event 同步：精确控制时间点
	60	+aclrtEvent event1, event2;
	61	+aclrtCreateEvent(&event1);
	62	+aclrtCreateEvent(&event2);
	63	+
	64	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_HOST_TO_DEVICE, stream1);
	65	+aclrtRecordEvent(event1, stream1); // 记录时间点1
	66	+
	67	+aclrtStreamWaitEvent(stream2, event1); // stream2 等待 event1
	68	+myKernel<<<..., stream2>>>();
	69	+aclrtRecordEvent(event2, stream2); // 记录时间点2
	70	+
	71	+aclrtSynchronizeEvent(event2); // 等待 event2 时间点完成
	72	+```
	73	+
	74	+### 原因2：错误选择同步方式
	75	+
	76	+解决方法：
	77	+- 单 Stream 等待：使用 aclrtSynchronizeStream，简单直接
	78	+- 跨 Stream 协调：使用 Event，stream1 Record，stream2 Wait
	79	+- 精细时间点控制：使用 Event，记录和等待特定时间点
	80	+- 避免过度设计：单流场景不要用 Event 同步
	81	+
	82	+场景选择示例：
	83	+```c
	84	+// 场景1：单 Stream 等待整批任务 → Stream 同步
	85	+aclrtStream stream;
	86	+aclrtCreateStream(&stream);
	87	+
	88	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	89	+myKernel<<<..., stream>>>();
	90	+aclrtMemcpyAsync(host, size, dst, size, ACL_MEMCPY_DEVICE_TO_HOST, stream);
	91	+
	92	+// 推荐：Stream 同步
	93	+aclrtSynchronizeStream(stream); // 等待整流完成
	94	+
	95	+// 不推荐：Event 同步（过度设计）
	96	+aclrtEvent event;
	97	+aclrtRecordEvent(event, stream);
	98	+aclrtSynchronizeEvent(event); // 多了一步 RecordEvent，复杂度高
	99	+
	100	+// 场景2：跨 Stream 协调 → Event 同步
	101	+aclrtStream stream1, stream2;
	102	+aclrtCreateStream(&stream1);
	103	+aclrtCreateStream(&stream2);
	104	+
	105	+aclrtMemcpyAsync(dst1, size, src1, size, ACL_MEMCPY_HOST_TO_DEVICE, stream1);
	106	+
	107	+// 推荐：Event 同步
	108	+aclrtEvent event;
	109	+aclrtRecordEvent(event, stream1); // stream1 记录时间点
	110	+aclrtStreamWaitEvent(stream2, event); // stream2 等待该时间点
	111	+
	112	+aclrtMemcpyAsync(dst2, size, src2, size, ACL_MEMCPY_HOST_TO_DEVICE, stream2);
	113	+
	114	+// 不推荐：Stream 同步（会阻塞 stream1，影响并发）
	115	+aclrtSynchronizeStream(stream1); // 阻塞 stream1 全部任务
	116	+aclrtMemcpyAsync(dst2, size, src2, size, ACL_MEMCPY_HOST_TO_DEVICE, stream2);
	117	+
	118	+// 场景3：精细时间点控制 → Event 同步
	119	+aclrtEvent event1, event2, event3;
	120	+aclrtCreateEvent(&event1);
	121	+aclrtCreateEvent(&event2);
	122	+aclrtCreateEvent(&event3);
	123	+
	124	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	125	+aclrtRecordEvent(event1, stream); // 记录复制完成时间点
	126	+
	127	+myKernel<<<..., stream>>>();
	128	+aclrtRecordEvent(event2, stream); // 记录算子完成时间点
	129	+
	130	+aclrtMemcpyAsync(host, size, dst, size, ACL_MEMCPY_DEVICE_TO_HOST, stream);
	131	+aclrtRecordEvent(event3, stream); // 记录输出复制完成时间点
	132	+
	133	+// 灵活等待特定时间点
	134	+aclrtSynchronizeEvent(event2); // 仅等待算子完成，不等待后续复制
	135	+```
	136	+
	137	+## 相关 issue
	138	+
	139	+- [Issue #477: aclrtSynchronizeStream报错507015](https://gitcode.com/cann/runtime/issues/477)

Adocs/04_FAQ/aclInit初始化失败常见原因排查.md+97-0

@@ -0,0 +1,97 @@
	1	+# aclInit初始化失败常见原因排查
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：调用aclInit接口返回参数错误码
	6	+
	7	+调用 aclInit 接口时返回 ACL_ERROR_RT_PARAM_INVALID（错误码 107000），表示参数无效。
	8	+
	9	+报错日志示例如下：
	10	+```
	11	+aclInit failed, ret = 107000
	12	+```
	13	+
	14	+现象2：调用aclInit接口返回重复初始化错误码
	15	+
	16	+调用 aclInit 接口时返回 ACL_ERROR_REPEAT_INITIALIZE（错误码 100002），表示重复初始化错误。
	17	+
	18	+报错日志示例如下：
	19	+```
	20	+aclInit failed, ret = 100002, repeat initialize
	21	+```
	22	+
	23	+现象3：配置文件解析失败
	24	+
	25	+aclInit 接口传入的配置文件路径错误或格式不正确，导致解析失败。
	26	+
	27	+报错日志示例如下：
	28	+```
	29	+Failed to parse config file: /path/to/acl.json
	30	+```
	31	+
	32	+## 可能原因
	33	+
	34	+1. 配置文件路径不存在或权限不足：aclInit 接口传入的配置文件路径错误，或文件不可读。
	35	+2. json 配置格式错误：配置文件括号层级超限（最多10层）、字段拼写错误、格式不符合 json 规范。
	36	+3. 多次 aclInit 调用配置不一致：重复调用 aclInit 时，配置必须保持一致，否则后续调用可能报错或配置无效。
	37	+
	38	+## 处理步骤
	39	+
	40	+### 原因1：配置文件路径不存在或权限不足
	41	+
	42	+解决方法：
	43	+- 使用绝对路径：传入完整的文件路径，例如 `/home/user/acl.json`
	44	+- 检查文件权限：确保文件可读，使用 `ls -l /path/to/acl.json` 查看权限
	45	+- 如果默认配置满足需求：传入 nullptr 或配置空 json 串 `{}`
	46	+
	47	+示例代码：
	48	+```c
	49	+// 使用绝对路径
	50	+aclError ret = aclInit("/home/user/config/acl.json");
	51	+
	52	+// 使用默认配置（传入 nullptr）
	53	+aclError ret = aclInit(nullptr);
	54	+
	55	+// 使用空配置（传入空 json）
	56	+// acl.json 内容：{}
	57	+aclError ret = aclInit("../acl.json");
	58	+```
	59	+
	60	+### 原因2：json 配置格式错误
	61	+
	62	+解决方法：
	63	+- 检查括号层级：json 文件内 `{` 层级最多10层，`[` 层级最多10层
	64	+- 验证字段拼写：参考 API 文档中的配置示例
	65	+- 使用 json 验证工具：在线 json 验证器检查格式正确性
	66	+
	67	+典型配置示例：
	68	+```json
	69	+{
	70	+ "defaultDevice":{
	71	+ "default_device":"0"
	72	+ }
	73	+}
	74	+```
	75	+
	76	+### 原因3：多次 aclInit 调用配置不一致
	77	+
	78	+解决方法：
	79	+- 保持配置一致：每次调用 aclInit 时使用相同的配置文件路径和内容
	80	+- 忽略重复初始化错误：aclInit 重复调用会返回 ACL_ERROR_REPEAT_INITIALIZE，可忽略该错误继续业务处理
	81	+- 成对调用初始化和去初始化：支持重复初始化和去初始化，时序为 `aclInit → 业务处理 → aclFinalize → aclInit → 业务处理 → aclFinalize`
	82	+
	83	+示例代码：
	84	+```c
	85	+// 正确的重复初始化方式
	86	+aclError ret1 = aclInit(nullptr); // 首次初始化
	87	+// ... 业务处理 ...
	88	+aclError ret2 = aclFinalize(); // 去初始化
	89	+
	90	+aclError ret3 = aclInit(nullptr); // 再次初始化（配置一致）
	91	+// ... 业务处理 ...
	92	+aclError ret4 = aclFinalize(); // 再次去初始化
	93	+```
	94	+
	95	+## 相关 issue
	96	+
	97	+暂无相关Issue。

Adocs/04_FAQ/aclrtMalloc内存申请失败常见原因.md+83-0

@@ -0,0 +1,83 @@
	1	+# aclrtMalloc内存申请失败常见原因
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：调用aclrtMalloc接口申请Device内存时返回内存申请失败错误码
	6	+
	7	+调用 aclrtMalloc 接口申请 Device 内存时返回错误码 207001（ACL_ERROR_RT_MEMORY_ALLOCATION），表示内存申请失败。
	8	+
	9	+报错日志示例如下：
	10	+```
	11	+aclrtMalloc failed, ret = 207001, size=10485760
	12	+[ERROR] RUNTIME: Failed to allocate device memory, device_id=0, size=10485760
	13	+```
	14	+
	15	+## 可能原因
	16	+
	17	+1. Device 内存不足或申请大小不合理：设备内存已被大量占用、申请大小超过设备内存总容量、并发申请超过系统资源上限，剩余容量不足以满足申请需求。
	18	+2. 内存分配策略配置不当：强制使用大页策略（如 ACL_MEM_MALLOC_HUGE_ONLY）但大页内存不足。
	19	+
	20	+## 处理步骤
	21	+
	22	+### 原因1：Device 内存不足或申请大小不合理
	23	+
	24	+解决方法：
	25	+- 检查内存容量：使用 npu-smi 工具查看设备内存使用情况
	26	+- 释放已分配内存：调用 aclrtFree 释放不再使用的内存
	27	+- 减少申请大小：优化业务逻辑，减少一次性申请的内存量
	28	+- 检查 size 参数：确认申请大小合理，不超过设备内存总容量
	29	+- 预分配并复用：初始化阶段预分配固定内存池，业务中复用
	30	+
	31	+命令示例：
	32	+```bash
	33	+# 查看设备内存使用情况
	34	+npu-smi info -t memory
	35	+
	36	+# 输出示例：
	37	+# Device ID: 0
	38	+# HBM Total: 32GB
	39	+# HBM Used: 28GB
	40	+# HBM Free: 4GB
	41	+```
	42	+
	43	+代码示例：
	44	+```c
	45	+// 检查内存使用情况后合理申请
	46	+void* devPtr = nullptr;
	47	+size_t size = 1024 * 1024; // 1MB
	48	+
	49	+// 先释放不再使用的内存
	50	+if (oldDevPtr != nullptr) {
	51	+ aclrtFree(oldDevPtr);
	52	+ oldDevPtr = nullptr;
	53	+}
	54	+
	55	+// 再申请新内存
	56	+aclError ret = aclrtMalloc(&devPtr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	57	+```
	58	+
	59	+### 原因2：内存分配策略配置不当
	60	+
	61	+解决方法：
	62	+- 理解策略含义：
	63	+ - ACL_MEM_MALLOC_HUGE_FIRST：优先大页，不足时退回普通页（推荐）
	64	+ - ACL_MEM_MALLOC_HUGE_ONLY：仅大页，不足时报错
	65	+ - ACL_MEM_MALLOC_NORMAL_ONLY：仅普通页
	66	+- 选择合适策略：根据实际内存容量和性能需求选择
	67	+- 调整策略为 HUGE_FIRST：避免强制大页导致的申请失败
	68	+
	69	+策略选择建议：
	70	+```c
	71	+// 推荐策略：大页优先，不足时退回普通页
	72	+aclError ret = aclrtMalloc(&devPtr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	73	+
	74	+// 不推荐策略：仅大页，大页不足时会失败
	75	+aclError ret = aclrtMalloc(&devPtr, size, ACL_MEM_MALLOC_HUGE_ONLY); // 可能失败
	76	+
	77	+// 普通页策略：适合小块内存或内存紧张场景
	78	+aclError ret = aclrtMalloc(&devPtr, size, ACL_MEM_MALLOC_NORMAL_ONLY);
	79	+```
	80	+
	81	+## 相关 issue
	82	+
	83	+- [Issue #476: aclrtMallocPhysical申请大页内存限制](https://gitcode.com/cann/runtime/issues/476)

Adocs/04_FAQ/aclrtMemcpyAsync在错误的Stream上下发失败.md+85-0

@@ -0,0 +1,85 @@
	1	+# aclrtMemcpyAsync在错误的Stream上下发失败
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：调用aclrtMemcpyAsync接口返回设备不匹配错误码
	6	+
	7	+调用 aclrtMemcpyAsync 异步内存复制接口时返回错误码 107003（ACL_ERROR_RT_STREAM_CONTEXT），表示Stream不在当前Device的Context中。
	8	+
	9	+典型错误场景：
	10	+```c
	11	+aclrtSetDevice(0); // 指定 Device 0
	12	+aclrtStream s0;
	13	+aclrtCreateStream(&s0); // 在 Device 0 上创建 Stream s0
	14	+
	15	+aclrtSetDevice(1); // 指定 Device 1
	16	+aclrtStream s1;
	17	+aclrtCreateStream(&s1); // 在 Device 1 上创建 Stream s1
	18	+
	19	+// 错误：在 Device 1 上通过 Device 0 的 Stream s0 下发任务
	20	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_DEVICE_TO_DEVICE, s0); // 失败
	21	+```
	22	+
	23	+报错日志示例如下：
	24	+```
	25	+aclrtMemcpyAsync failed, ret = 107003, stream not in current context
	26	+[ERROR] RUNTIME: Stream not in current context, stream device mismatch
	27	+```
	28	+
	29	+## 可能原因
	30	+
	31	+1. Stream 与当前 Device 不属于同一 Device：Stream 在创建时绑定到特定 Device，不能在另一个 Device 上下发任务。
	32	+
	33	+## 处理步骤
	34	+
	35	+### 原因1：Stream 与当前 Device 不属于同一 Device
	36	+
	37	+解决方法：
	38	+- 理解 Stream 归属概念：Stream 属于创建时的 Device，与 Device 绑定
	39	+- 在正确 Device 上下发：切换 Device 后使用属于该 Device 的 Stream
	40	+- 每个 Device 创建独立 Stream：避免跨 Device 混用 Stream
	41	+
	42	+正确使用示例：
	43	+```c
	44	+// Device 0 的操作
	45	+aclrtSetDevice(0);
	46	+aclrtStream s0;
	47	+aclrtCreateStream(&s0); // s0 属于 Device 0
	48	+
	49	+// Device 0 上使用 s0 下发任务
	50	+aclrtMemcpyAsync(dst0, size, src0, size, ACL_MEMCPY_DEVICE_TO_DEVICE, s0); // 正确
	51	+
	52	+// Device 1 的操作
	53	+aclrtSetDevice(1);
	54	+aclrtStream s1;
	55	+aclrtCreateStream(&s1); // s1 属于 Device 1
	56	+
	57	+// Device 1 上使用 s1 下发任务
	58	+aclrtMemcpyAsync(dst1, size, src1, size, ACL_MEMCPY_DEVICE_TO_DEVICE, s1); // 正确
	59	+
	60	+// 错误示例：在 Device 1 上使用 Device 0 的 Stream s0
	61	+aclrtMemcpyAsync(dst1, size, src1, size, ACL_MEMCPY_DEVICE_TO_DEVICE, s0); // 失败！
	62	+```
	63	+
	64	+多 Device 编程最佳实践：
	65	+```c
	66	+// 为每个 Device 创建独立的 Stream
	67	+aclrtStream deviceStreams[2];
	68	+
	69	+aclrtSetDevice(0);
	70	+aclrtCreateStream(&deviceStreams[0]); // Device 0 的 Stream
	71	+
	72	+aclrtSetDevice(1);
	73	+aclrtCreateStream(&deviceStreams[1]); // Device 1 的 Stream
	74	+
	75	+// 切换 Device 后使用对应的 Stream
	76	+aclrtSetDevice(0);
	77	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_DEVICE_TO_DEVICE, deviceStreams[0]); // 正确
	78	+
	79	+aclrtSetDevice(1);
	80	+aclrtMemcpyAsync(dst, size, src, size, ACL_MEMCPY_DEVICE_TO_DEVICE, deviceStreams[1]); // 正确
	81	+```
	82	+
	83	+## 相关 issue
	84	+
	85	+暂无相关Issue。

Adocs/04_FAQ/aclrtSetDevice调用失败.md+66-0

@@ -0,0 +1,66 @@
	1	+# aclrtSetDevice调用失败
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：调用aclrtSetDevice接口返回设备ID无效错误码
	6	+
	7	+调用 aclrtSetDevice 接口失败，返回 ACL_ERROR_RT_INVALID_DEVICEID（错误码 107001），表示设备ID无效。
	8	+
	9	+报错日志示例如下：
	10	+```
	11	+aclrtSetDevice failed, ret = 107001, deviceId=5
	12	+```
	13	+
	14	+现象2：调用aclrtSetDevice接口失败且驱动未正确加载
	15	+
	16	+当 NPU 驱动未安装或未正确加载时，调用 aclrtSetDevice 接口也会失败，可能返回 ACL_ERROR_RT_INVALID_DEVICEID（错误码 107001）。
	17	+
	18	+报错日志示例如下：
	19	+```
	20	+aclrtSetDevice failed, ret = 107001, deviceId=0
	21	+[ERROR] RUNTIME: Failed to get phy dev id by logic dev id, driver may not be loaded
	22	+```
	23	+
	24	+## 可能原因
	25	+
	26	+1. Device ID 超出可用范围：指定的设备ID超过了系统实际可用的设备数量。
	27	+2. 驱动未正确加载：NPU驱动未安装或未正确加载到系统。
	28	+
	29	+## 处理步骤
	30	+
	31	+### 原因1：Device ID 超出可用范围
	32	+
	33	+解决方法：
	34	+- 查询可用设备数量：调用 aclrtGetDeviceCount 获取系统中的设备总数
	35	+- 使用正确的设备ID：设备ID从0开始编号，范围为 `[0, deviceCount-1]`
	36	+
	37	+示例代码：
	38	+```c
	39	+uint32_t deviceCount = 0;
	40	+aclError ret = aclrtGetDeviceCount(&deviceCount);
	41	+if (ret == ACL_SUCCESS) {
	42	+ // 设备ID范围：0 到 deviceCount-1
	43	+ int32_t deviceId = 0; // 使用有效的设备ID
	44	+ ret = aclrtSetDevice(deviceId);
	45	+}
	46	+```
	47	+
	48	+### 原因2：驱动未正确加载
	49	+
	50	+解决方法：
	51	+- 检查驱动安装：确认 CANN 驱动已正确安装
	52	+- 查看驱动状态：使用系统命令检查驱动加载情况
	53	+- 重装驱动：如驱动未加载，参考昇腾社区文档重新安装驱动
	54	+
	55	+命令示例：
	56	+```bash
	57	+# 查看驱动版本
	58	+cat /usr/local/Ascend/version.info
	59	+
	60	+# 查看驱动加载状态
	61	+lsmod \| grep ascend
	62	+```
	63	+
	64	+## 相关 issue
	65	+
	66	+暂无相关Issue。

Ddocs/04_FAQ/低版本内核使用asan导致算子执行失败.md+0-35

	@@ -1,35 +0,0 @@
1	-# 低版本内核使用asan导致算子执行失败
2	-
3	-## 问题现象描述
4	-
5	-执行算子时，算子输入数据正确，但输出数据异常，全为0，Host侧plog日志中的报错示例如下：
6	-
7	-```
8	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.036.721 [stars_engine.cc:1321]2082291 ProcLogicCqReport:[INIT][DEFAULT]Task run failed, device_id=0, stream_id=2, task_id=1, sqe_type=0(ffts), errType=0x1(task exception), sqSwStatus=0
9	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.049.079 [device_error_proc.cc:1218]2082291 ProcessStarsCoreErrorInfo:[INIT][DEFAULT]report error module_type=5, module_name=EZ9999
10	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.049.115 [device_error_proc.cc:1218]2082291 ProcessStarsCoreErrorInfo:[INIT][DEFAULT]The error from device(chipId:3, dieId:0), serial number is 20, there is an aivec error exception, core id is 4, error code = 0, dump info: pc start: 0x12c0c001406c, current: 0x12c0c00140fc, vec error info: 0x600ed4063d, mte error info: 0x8d0600008c, ifu error info: 0x70f016e068500, ccu error info: 0x28000037, cube error info: 0, biu error info: 0, aic error mask: 0x6500020bd00028c, para base: 0x12c0803e5000.
11	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.049.300 [device_error_proc.cc:1230]2082291 ProcessStarsCoreErrorInfo:[INIT][DEFAULT]report error module_type=5, module_name=EZ9999
12	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.049.321 [device_error_proc.cc:1230]2082291 ProcessStarsCoreErrorInfo:[INIT][DEFAULT]The extend info: errcode:(0, 0x200000000000000, 0) errorStr: The MPU address access is invalid. fixp_error0 info: 0x600008c, fixp_error1 info: 0x8d fsmId:1, tslot:3, thread:0, ctxid:0, blk:0, sublk:0, subErrType:4.
13	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.049.519 [stream.cc:3084]2082291 EnterFailureAbort:[INIT][DEFAULT]stream_id=2 enter failure abort.
14	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.049.558 [davinic_kernel_task.cc:1321]2082291 SetStarsResultForDavinciTask:[INIT][DEFAULT]AIV Kernel happen error, retCode=0x31.
15	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.340 [davinic_kernel_task.cc:1219]2082291 PreCheckTaskErr:[INIT][DEFAULT]report error module_type=5, module_name=EZ9999
16	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.365 [davinic_kernel_task.cc:1219]2082291 PreCheckTaskErr:[INIT][DEFAULT]Kernel task happen error, retCode=0x31, [vector core exception].
17	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.474 [stream.cc:1079]2082291 GetError:[INIT][DEFAULT]Stream Synchronize failed, stream_id=2, retCode=0x31, [vector core exception].
18	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.496 [stream.cc:1082]2082291 GetError:[INIT][DEFAULT]report error module_type=5, module_name=EZ9999
19	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.517 [stream.cc:1082]2082291 GetError:[INIT][DEFAULT]AIV Kernel happen error, retCode=0x31.
20	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.941 [davinic_kernel_task.cc:1143]2082291 PrintErrorInfoForDavinciTask:[INIT][DEFAULT]Aicore kernel execute failed, device_id=0, stream_id=2, report_stream_id=2, task_id=1, flip_num=0, fault kernel_name=Add_ee98c6628030785f610b924ab1557b31_high_performance_210000000, fault kernel info ext=none, program id=0, hash=3838710036602041089.
21	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.051.013 [davinic_kernel_task.cc:1082]2082291 GetArgsInfo:[INIT][DEFAULT][AIC_INFO] args(0 to 9) after execute:0, 0, 0, 0, 0, 0, 0, 0, 0,
22	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.051.046 [davinic_kernel_task.cc:1085]2082291 GetArgsInfo:[INIT][DEFAULT]tilingKey = 210000000, print 1 Times totalLen=(9*8)Bytes, argsSize=72, blockDim=1
23	-[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.051.088 [davinic_kernel_task.cc:1147]2082291 PrintErrorInfoForDavinciTask:[INIT][DEFAULT][AIC_INFO] after execute:args print end
24	-```
25	-
26	-## 可能原因
27	-
28	-用户程序编译选项中启动了地址消毒（-lasan），但低版本内核（5.10以下版本，不含5.10）不支持使用asan工具，导致执行算子时拷贝输出数据异常。
29	-
30	-可使用uname -r命令查看内核版本。
31	-
32	-## 处理步骤
33	-
34	-- 解决方法1：升级内核版本到5.10或更高版本。
35	-- 解决方法2：在用户程序编译选项中去掉地址消毒（-lasan）。

Adocs/04_FAQ/多Device场景下Stream跨Device下发失败.md+104-0

@@ -0,0 +1,104 @@
	1	+# 多Device场景下Stream跨Device下发失败
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：在非所属Device的Stream上下发算子失败
	6	+
	7	+在多Device场景下，当尝试在与当前Device无所属关系的Stream上下发算子时，任务下发失败。
	8	+
	9	+错误代码示例：
	10	+```cpp
	11	+aclrtSetDevice(0);
	12	+aclrtStream s0;
	13	+aclrtCreateStream(&s0);
	14	+myKernel<<<8, nullptr, s0>>>(); // ✓ Device 0上创建，Device 0上使用
	15	+
	16	+aclrtSetDevice(1);
	17	+aclrtStream s1;
	18	+aclrtCreateStream(&s1);
	19	+myKernel<<<8, nullptr, s1>>>(); // ✓ Device 1上创建，Device 1上使用
	20	+
	21	+myKernel<<<8, nullptr, s0>>>(); // ✗ 在Device 1上通过Stream s0下发算子，失败
	22	+```
	23	+
	24	+现象2：在非所属Device的Stream上调用异步复制接口失败
	25	+
	26	+当Stream所属的Device和当前操作的Device不相同时，在此Stream上调用aclrtMemcpyAsync会失败。
	27	+
	28	+现象3：Event和Stream关联到不同Device时操作失败
	29	+
	30	+当Event和Stream关联到不同的Device上时，调用aclrtRecordEvent或aclrtStreamWaitEvent会失败。
	31	+
	32	+## 可能原因
	33	+
	34	+1. Stream具有Device所属关系，Stream在其所属的Device上创建后，只能在该Device上使用。
	35	+2. Event也具有Device所属关系，Event只能关联到与其所属Device相同的Stream上。
	36	+3. 当前操作的Device与Stream/Event所属Device不匹配，违反了资源所属关系约束。
	37	+
	38	+## 处理步骤
	39	+
	40	+方案一：在正确的Device上使用Stream
	41	+
	42	+确保在Stream所属的Device上进行操作：
	43	+
	44	+```cpp
	45	+// 正确示例：在Stream所属Device上使用
	46	+aclrtSetDevice(0);
	47	+aclrtStream s0;
	48	+aclrtCreateStream(&s0);
	49	+myKernel<<<8, nullptr, s0>>>(); // ✓ 正确：Device 0上创建，Device 0上使用
	50	+
	51	+aclrtSetDevice(1);
	52	+aclrtStream s1;
	53	+aclrtCreateStream(&s1);
	54	+myKernel<<<8, nullptr, s1>>>(); // ✓ 正确：Device 1上创建，Device 1上使用
	55	+```
	56	+
	57	+方案二：使用Event实现跨Device同步
	58	+
	59	+如果需要跨Device协调任务，使用Event进行同步而不是跨Device使用Stream：
	60	+
	61	+```cpp
	62	+// 在Device 0上创建Event并记录
	63	+aclrtSetDevice(0);
	64	+aclrtStream s0;
	65	+aclrtCreateStream(&s0);
	66	+aclrtEvent event;
	67	+aclrtCreateEvent(&event);
	68	+aclrtRecordEvent(event, s0);
	69	+
	70	+// 切换到Device 1，等待Device 0的事件
	71	+aclrtSetDevice(1);
	72	+aclrtStream s1;
	73	+aclrtCreateStream(&s1);
	74	+aclrtStreamWaitEvent(s1, event); // Device 1上的Stream等待Device 0的事件
	75	+myKernel<<<8, nullptr, s1>>>();
	76	+```
	77	+
	78	+方案三：检查Event和Stream的Device所属关系
	79	+
	80	+确保Event和Stream关联到相同的Device：
	81	+
	82	+```cpp
	83	+// 错误示例
	84	+aclrtSetDevice(0);
	85	+aclrtEvent event;
	86	+aclrtCreateEvent(&event); // Event属于Device 0
	87	+
	88	+aclrtSetDevice(1);
	89	+aclrtStream s1;
	90	+aclrtCreateStream(&s1); // Stream属于Device 1
	91	+aclrtRecordEvent(event, s1); // ✗ 错误：Event和Stream属于不同Device
	92	+
	93	+// 正确示例
	94	+aclrtSetDevice(0);
	95	+aclrtEvent event;
	96	+aclrtCreateEvent(&event); // Event属于Device 0
	97	+aclrtStream s0;
	98	+aclrtCreateStream(&s0); // Stream属于Device 0
	99	+aclrtRecordEvent(event, s0); // ✓ 正确：Event和Stream属于同一Device
	100	+```
	101	+
	102	+## 相关 issue
	103	+
	104	+- [Issue #344: device、context、stream关系咨询](https://gitcode.com/cann/runtime/issues/344)

Adocs/04_FAQ/如何理解默认Device和默认Stream机制.md+117-0

@@ -0,0 +1,117 @@
	1	+# 如何理解默认Device和默认Stream机制
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：不调用aclrtSetDevice直接使用其他接口失败
	6	+
	7	+未调用 aclrtSetDevice 配置设备，直接调用 aclrtMalloc 等接口时失败。
	8	+
	9	+错误代码示例：
	10	+```c
	11	+aclInit(nullptr);
	12	+aclrtMalloc(&devPtr, size, ACL_MEM_MALLOC_HUGE_FIRST); // 失败：没有指定Device
	13	+```
	14	+
	15	+现象2：对默认Stream的创建时机和使用方式不理解
	16	+
	17	+不清楚默认Stream何时自动创建和销毁，以及与显式创建Stream的区别。
	18	+
	19	+典型困惑：
	20	+```c
	21	+aclrtSetDevice(0); // 自动创建默认Stream
	22	+aclrtMemcpyAsync(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE, nullptr); // nullptr是默认Stream？
	23	+```
	24	+
	25	+现象3：混淆显式创建Stream和默认Stream的使用场景
	26	+
	27	+不理解何时传 nullptr（默认Stream）、何时传显式创建的Stream对象。
	28	+
	29	+## 可能原因
	30	+
	31	+1. 未配置 aclInit 的 defaultDevice 功能：aclInit 默认不设置 defaultDevice，需要显式调用 aclrtSetDevice。
	32	+2. 混淆默认Stream和显式创建Stream的使用场景：不理解默认Stream的自动创建机制和销毁时机。
	33	+3. 不理解 Stream 传参规则：需要Stream参数的接口，默认Stream传 nullptr，显式创建Stream传实际对象。
	34	+
	35	+## 处理步骤
	36	+
	37	+### 原因1：未配置 aclInit 的 defaultDevice 功能
	38	+
	39	+解决方法：
	40	+- 配置 defaultDevice：在 aclInit 的 json 配置文件中设置 defaultDevice
	41	+- 理解 defaultDevice 作用：启用后可不显式调用 aclrtSetDevice，接口内部自动进行隐式 aclrtSetDevice
	42	+
	43	+配置示例：
	44	+```json
	45	+{
	46	+ "defaultDevice":{
	47	+ "default_device":"0"
	48	+ }
	49	+}
	50	+```
	51	+
	52	+使用示例：
	53	+```c
	54	+// 启用 defaultDevice 后
	55	+aclError ret = aclInit("../acl.json"); // json中配置了defaultDevice=0
	56	+
	57	+// 可以直接调用运行时接口，无需显式 aclrtSetDevice
	58	+aclrtMalloc(&devPtr, size, ACL_MEM_MALLOC_HUGE_FIRST); // 成功
	59	+
	60	+// 去初始化
	61	+aclrtResetDeviceForce(0);
	62	+aclFinalize();
	63	+```
	64	+
	65	+### 原因2：混淆默认Stream和显式创建Stream的使用场景
	66	+
	67	+解决方法：
	68	+- 理解默认Stream创建时机：aclrtSetDevice 或 aclrtCreateContext 时自动创建默认Stream
	69	+- 理解默认Stream销毁时机：aclrtResetDevice 或 aclrtResetDeviceForce 时自动销毁默认Stream
	70	+- 区分两种Stream：显式创建的Stream需调用 aclrtDestroyStream 销毁，默认Stream不能显式销毁
	71	+
	72	+默认Stream使用示例：
	73	+```c
	74	+aclrtSetDevice(0); // 自动创建默认Stream
	75	+
	76	+// 默认Stream传 nullptr
	77	+aclrtMemcpyAsync(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE, nullptr);
	78	+
	79	+// 同步默认Stream
	80	+aclrtSynchronizeStream(nullptr); // nullptr 表示默认Stream
	81	+
	82	+aclrtResetDevice(0); // 自动销毁默认Stream
	83	+```
	84	+
	85	+显式创建Stream示例：
	86	+```c
	87	+aclrtSetDevice(0);
	88	+aclrtStream stream;
	89	+aclrtCreateStream(&stream); // 显式创建Stream
	90	+
	91	+// 显式Stream传实际对象
	92	+aclrtMemcpyAsync(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	93	+
	94	+// 销毁显式创建的Stream
	95	+aclrtDestroyStream(stream); // 必须显式销毁
	96	+aclrtResetDevice(0);
	97	+```
	98	+
	99	+### 原因3：不理解 Stream 传参规则
	100	+
	101	+解决方法：
	102	+- 需要Stream参数的接口（如 aclrtMemcpyAsync）：默认Stream传 nullptr，显式创建Stream传实际对象
	103	+- 不需要Stream参数的接口（如 aclrtMemcpy）：不使用默认Stream，属于同步接口
	104	+
	105	+传参规则总结：
	106	+```c
	107	+// 异步接口：需要Stream参数
	108	+aclrtMemcpyAsync(..., nullptr); // 使用默认Stream
	109	+aclrtMemcpyAsync(..., stream); // 使用显式创建的Stream
	110	+
	111	+// 同步接口：不需要Stream参数
	112	+aclrtMemcpy(..., ACL_MEMCPY_HOST_TO_DEVICE); // 不使用Stream，同步执行
	113	+```
	114	+
	115	+## 相关 issue
	116	+
	117	+- [Issue #344: device、context、stream关系咨询](https://gitcode.com/cann/runtime/issues/344)

Adocs/04_FAQ/如何获取和解读Runtime异步错误码.md+169-0

@@ -0,0 +1,169 @@
	1	+# 如何获取和解读Runtime异步错误码
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：异步接口返回成功但Device侧实际执行失败
	6	+
	7	+调用Runtime异步接口（如aclrtMemcpyAsync、Kernel Launch等）时，接口返回成功（ACL_RT_SUCCESS），但实际执行时Device侧发生了错误。
	8	+
	9	+典型场景：
	10	+```cpp
	11	+aclrtStream stream;
	12	+aclrtCreateStream(&stream);
	13	+
	14	+// 异步接口返回成功，仅表示任务下发成功
	15	+aclError error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	16	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	17	+// error == ACL_RT_SUCCESS，但Device侧可能执行失败
	18	+```
	19	+
	20	+现象2：异步错误延迟传播导致难以定位
	21	+
	22	+Device侧的异步错误会延迟到后续某个Runtime接口调用时返回，可能不是实际发生错误的接口，导致错误定位困难。
	23	+
	24	+报错日志示例如下：
	25	+```
	26	+[ERROR] RUNTIME: Aicore kernel execute failed, ret=507015, aicore exception
	27	+fault kernel_name=Add_ee98c6628030785f610b924ab1557b31
	28	+```
	29	+
	30	+现象3：遇错继续模式下多个错误覆盖
	31	+
	32	+在遇错继续模式下，如果Stream上多个任务执行失败，后发生的错误可能覆盖先前的错误信息，导致无法获取首次错误。
	33	+
	34	+// 后续操作基于错误假设继续执行
	35	+myKernel<<<8, nullptr, stream>>>();
	36	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	37	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	38	+aclrtSynchronizeStream(stream);
	39	+// 此时才发现错误，但难以定位是哪个异步任务失败
	40	+
	41	+## 可能原因
	42	+
	43	+1. 异步执行机制：Runtime异步接口在Host侧下发任务后立即返回，Device侧任务异步执行，接口返回值仅反映Host侧参数校验和任务下发是否成功，无法报告Device侧的实际执行错误。
	44	+2. 错误传播延迟：Device侧的异步错误会延迟到后续某个Runtime接口调用时返回，可能不是实际发生错误的接口，导致错误定位困难。
	45	+3. 错误覆盖问题：在遇错继续模式下，如果Stream上多个任务执行失败，后发生的错误可能覆盖先前的错误信息，导致无法获取首次错误。
	46	+
	47	+典型异步错误示例（通过 aclrtSynchronizeStream 获取）：
	48	+```
	49	+[ERROR] RUNTIME: Aicore kernel execute failed, ret=507015, aicore exception
	50	+fault kernel_name=Add_ee98c6628030785f610b924ab1557b31
	51	+```
	52	+
	53	+## 处理步骤
	54	+
	55	+### 方法1：立即同步获取异步错误
	56	+
	57	+在异步接口调用后立即调用同步接口，获取Device侧的实际执行结果：
	58	+```cpp
	59	+aclrtStream stream;
	60	+aclrtCreateStream(&stream);
	61	+
	62	+// 下发异步任务
	63	+aclError error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	64	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	65	+
	66	+// 立即同步并获取错误码
	67	+error = aclrtSynchronizeStream(stream);
	68	+if (error != ACL_RT_SUCCESS) {
	69	+ // 处理错误
	70	+ printf("Async task failed with error: %d\n", error);
	71	+ char *errMsg = aclGetRecentErrMsg();
	72	+ printf("Error message: %s\n", errMsg);
	73	+}
	74	+
	75	+aclrtDestroyStream(stream);
	76	+```
	77	+
	78	+### 方法2：使用错误查询接口
	79	+
	80	+使用`aclrtPeekAtLastError`或`aclrtGetLastError`查询当前线程的错误状态：
	81	+
	82	+```cpp
	83	+aclrtStream stream;
	84	+aclrtCreateStream(&stream);
	85	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	86	+
	87	+// 下发多个异步任务
	88	+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	89	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	90	+myKernel<<<8, nullptr, stream>>>();
	91	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	92	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	93	+
	94	+// 同步并检查错误
	95	+aclrtSynchronizeStream(stream);
	96	+
	97	+// 查看错误但不重置
	98	+aclError lastError = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL);
	99	+if (lastError != ACL_RT_SUCCESS) {
	100	+ printf("Last error: %d\n", lastError);
	101	+}
	102	+
	103	+// 获取并重置错误状态
	104	+lastError = aclrtGetLastError(ACL_RT_THREAD_LEVEL);
	105	+// 此时错误状态已重置为ACL_RT_SUCCESS
	106	+```
	107	+
	108	+### 方法3：获取详细错误信息
	109	+
	110	+通过`aclGetRecentErrMsg()`获取详细的错误描述信息：
	111	+
	112	+```cpp
	113	+aclError error = aclrtSynchronizeDevice();
	114	+if (error != ACL_RT_SUCCESS) {
	115	+ char *errMsg = aclGetRecentErrMsg();
	116	+ if (errMsg != nullptr) {
	117	+ printf("Error detail: %s\n", errMsg);
	118	+ }
	119	+
	120	+ // 根据错误码进行分类处理
	121	+ // 实际错误码定义请参考 rt_error_codes.h 头文件
	122	+ switch (error) {
	123	+ case ACL_ERROR_RT_PARAM_INVALID: // 107000
	124	+ printf("Invalid parameter\n");
	125	+ break;
	126	+ case ACL_ERROR_RT_INVALID_DEVICEID: // 107001
	127	+ printf("Invalid device ID\n");
	128	+ break;
	129	+ // ... 其他错误码处理
	130	+ }
	131	+}
	132	+```
	133	+
	134	+### 方法4：结合遇错即停模式
	135	+
	136	+配置遇错即停模式，在首个错误发生时停止执行，避免错误传播：
	137	+```cpp
	138	+aclrtStream stream;
	139	+aclrtCreateStream(&stream);
	140	+
	141	+// 配置遇错即停
	142	+aclError error = aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	143	+if (error != ACL_RT_SUCCESS) {
	144	+ printf("Failed to set failure mode\n");
	145	+ return error;
	146	+}
	147	+
	148	+// 下发任务
	149	+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	150	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	151	+myKernel<<<8, nullptr, stream>>>();
	152	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	153	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	154	+
	155	+// 同步并检查
	156	+error = aclrtSynchronizeStream(stream);
	157	+if (error != ACL_RT_SUCCESS) {
	158	+ // 首个错误发生时即停止，便于定位
	159	+ printf("First error occurred: %d\n", error);
	160	+ char *errMsg = aclGetRecentErrMsg();
	161	+ printf("Error message: %s\n", errMsg);
	162	+}
	163	+
	164	+aclrtDestroyStream(stream);
	165	+```
	166	+
	167	+## 相关 issue
	168	+- [Issue #477: aclrtSynchronizeStream报错507015及plog日志分析](https://gitcode.com/cann/runtime/issues/477)
	169	+- [Issue #480: aicore异常报错分析](https://gitcode.com/cann/runtime/issues/480)

Adocs/04_FAQ/如何选择合适的内存分配策略.md+145-0

@@ -0,0 +1,145 @@
	1	+# 如何选择合适的内存分配策略
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：内存分配性能不佳
	6	+
	7	+内存分配速度慢、访问性能低，选择的页类型不适合业务场景。
	8	+
	9	+典型场景示例：
	10	+```c
	11	+// 小块内存却配置了HUGE_ONLY策略，实际占用远超需求
	12	+aclrtMalloc(&devPtr, 1024, ACL_MEM_MALLOC_HUGE_ONLY); // 实际占用 2MB（大页对齐）
	13	+```
	14	+
	15	+现象2：内存申请失败
	16	+
	17	+配置特定策略后内存申请失败，如强制使用大页但大页内存不足。
	18	+
	19	+现象3：P2P场景性能不佳
	20	+
	21	+跨Device复制场景使用普通策略而非P2P策略，性能不佳。
	22	+
	23	+## 可能原因
	24	+
	25	+1. 普通页 vs 大页选择不当：不理解普通页和大页的性能特性、容量占用差异。
	26	+2. P2P 场景未使用对应策略：跨 Device 复制场景使用普通策略，性能不佳。
	27	+3. 频繁申请释放内存：未预分配复用，导致性能损耗。
	28	+
	29	+## 处理步骤
	30	+
	31	+### 原因1：普通页 vs 大页选择不当
	32	+
	33	+解决方法：
	34	+- 理解页类型特性：
	35	+ - 普通页（4K）：适合小块、短生命周期、数量多的内存申请，减少对齐浪费
	36	+ - 2M大页：适合大块、长期驻留、频繁访问的内存，减少页表项、扩大TLB覆盖
	37	+ - 1G大页：适合超大块、长期驻留、性能敏感的内存（部分产品支持）
	38	+- 按数据大小选择策略：
	39	+ - ≤1M：使用 ACL_MEM_MALLOC_NORMAL_ONLY 或 HUGE_FIRST（自动降级为普通页）
	40	+ - >1M：使用 ACL_MEM_MALLOC_HUGE_FIRST（优先大页）
	41	+- 了解 HUGE_FIRST 机制：申请大小 ≤1M 时，即使配置 HUGE_FIRST 也使用普通页
	42	+
	43	+策略选择示例：
	44	+```c
	45	+// 小块内存（≤1M）：普通页或 HUGE_FIRST（自动降级）
	46	+size_t smallSize = 1024 * 100; // 100KB
	47	+aclrtMalloc(&ptr1, smallSize, ACL_MEM_MALLOC_NORMAL_ONLY); // 普通页，占用约100KB
	48	+aclrtMalloc(&ptr2, smallSize, ACL_MEM_MALLOC_HUGE_FIRST); // 自动使用普通页
	49	+
	50	+// 大块内存（>1M）：大页优先
	51	+size_t largeSize = 1024 * 1024 * 10; // 10MB
	52	+aclrtMalloc(&ptr3, largeSize, ACL_MEM_MALLOC_HUGE_FIRST); // 大页，占用约10MB（对齐到2M）
	53	+aclrtMalloc(&ptr4, largeSize, ACL_MEM_MALLOC_HUGE_ONLY); // 仅大页，不足时报错
	54	+
	55	+// 强制大页场景：对大页性能有强依赖
	56	+aclrtMalloc(&ptr5, largeSize, ACL_MEM_MALLOC_HUGE_ONLY); // 大页不足时报错，尽早暴露问题
	57	+```
	58	+
	59	+### 原因2：P2P 场景未使用对应策略
	60	+
	61	+解决方法：
	62	+- 理解 P2P 场景：跨 Device 数据复制，内存需支持跨 Device 访问
	63	+- 使用 P2P 策略：选择带 P2P 后缀的策略（HUGE_FIRST_P2P、HUGE_ONLY_P2P、NORMAL_ONLY_P2P）
	64	+- 使内存属性匹配访问路径：P2P 策略确保内存属性与跨 Device 访问匹配
	65	+
	66	+P2P 策略示例：
	67	+```c
	68	+// 跨 Device 复制场景
	69	+aclrtSetDevice(0);
	70	+aclrtDeviceEnablePeerAccess(1, 0); // 开启 Device 0→1 的 P2P
	71	+
	72	+aclrtSetDevice(1);
	73	+aclrtDeviceEnablePeerAccess(0, 0); // 开启 Device 1→0 的 P2P
	74	+
	75	+// Device 0 上申请用于 P2P 的内存
	76	+aclrtSetDevice(0);
	77	+void* dev0Mem = nullptr;
	78	+aclrtMalloc(&dev0Mem, size, ACL_MEM_MALLOC_HUGE_FIRST_P2P); // P2P 策略
	79	+
	80	+// Device 1 上申请用于 P2P 的内存
	81	+aclrtSetDevice(1);
	82	+void* dev1Mem = nullptr;
	83	+aclrtMalloc(&dev1Mem, size, ACL_MEM_MALLOC_HUGE_FIRST_P2P); // P2P 策略
	84	+
	85	+// 跨 Device 复制
	86	+aclrtMemcpy(dev1Mem, size, dev0Mem, size, ACL_MEMCPY_DEVICE_TO_DEVICE);
	87	+```
	88	+
	89	+### 原因3：频繁申请释放内存
	90	+
	91	+解决方法：
	92	+- 预分配内存池：初始化阶段预分配固定大小的内存
	93	+- 业务侧复用：避免频繁调用 aclrtMalloc/aclrtFree
	94	+- 减少性能损耗：预分配复用可显著提升性能
	95	+
	96	+内存池示例：
	97	+```c
	98	+// 初始化阶段预分配内存池
	99	+#define POOL_SIZE 10
	100	+#define MEMORY_BLOCK_SIZE (1024 * 1024) // 1MB
	101	+
	102	+void* memoryPool[POOL_SIZE];
	103	+bool memoryPoolUsed[POOL_SIZE] = {false};
	104	+
	105	+void InitMemoryPool() {
	106	+ for (int i = 0; i < POOL_SIZE; i++) {
	107	+ aclrtMalloc(&memoryPool[i], MEMORY_BLOCK_SIZE, ACL_MEM_MALLOC_HUGE_FIRST);
	108	+ memoryPoolUsed[i] = false;
	109	+ }
	110	+}
	111	+
	112	+// 业务中申请内存（从池中获取）
	113	+void* AllocateFromPool() {
	114	+ for (int i = 0; i < POOL_SIZE; i++) {
	115	+ if (!memoryPoolUsed[i]) {
	116	+ memoryPoolUsed[i] = true;
	117	+ return memoryPool[i];
	118	+ }
	119	+ }
	120	+ return nullptr; // 池已满
	121	+}
	122	+
	123	+// 业务中释放内存（归还到池）
	124	+void ReleaseToPool(void* ptr) {
	125	+ for (int i = 0; i < POOL_SIZE; i++) {
	126	+ if (memoryPool[i] == ptr) {
	127	+ memoryPoolUsed[i] = false;
	128	+ return;
	129	+ }
	130	+ }
	131	+}
	132	+
	133	+// 清理阶段释放池
	134	+void CleanupMemoryPool() {
	135	+ for (int i = 0; i < POOL_SIZE; i++) {
	136	+ if (memoryPool[i] != nullptr) {
	137	+ aclrtFree(memoryPool[i]);
	138	+ }
	139	+ }
	140	+}
	141	+```
	142	+
	143	+## 相关 issue
	144	+
	145	+- [Issue #476: aclrtMallocPhysical申请大页内存限制](https://gitcode.com/cann/runtime/issues/476)

Adocs/04_FAQ/如何通过plog日志定位Device侧异常.md+194-0

@@ -0,0 +1,194 @@
	1	+# 如何通过plog日志定位Device侧异常
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：算子执行失败但无法定位错误原因和位置
	6	+
	7	+Runtime应用在Device侧执行任务时发生异常，Host侧返回错误码，但无法定位具体的错误原因和位置。
	8	+
	9	+报错日志示例如下：
	10	+```
	11	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.036.721 [stars_engine.cc:1321]2082291 ProcLogicCqReport:[INIT][DEFAULT]Task run failed, device_id=0, stream_id=2, task_id=1, sqe_type=0(ffts), errType=0x1(task exception), sqSwStatus=0
	12	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.365 [davinic_kernel_task.cc:1219]2082291 PreCheckTaskErr:[INIT][DEFAULT]Kernel task happen error, retCode=0x31, [vector core exception].
	13	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.941 [davinic_kernel_task.cc:1143]2082291 PrintErrorInfoForDavinciTask:[INIT][DEFAULT]Aicore kernel execute failed, device_id=0, stream_id=2, report_stream_id=2, task_id=1, flip_num=0, fault kernel_name=Add_ee98c6628030785f610b924ab1557b31_high_performance_210000000
	14	+```
	15	+
	16	+现象2：Device侧任务异常退出或内存访问越界
	17	+
	18	+Device侧任务异常退出、内存访问越界或内核执行超时。
	19	+
	20	+现象3：需要通过plog日志进行深入分析
	21	+
	22	+常规错误码不足以定位问题，需要通过plog日志获取Device侧异常的详细信息。
	23	+
	24	+## 可能原因
	25	+
	26	+1. Device侧硬件异常：包括向量核异常（vector core exception）、立方核异常（cube core exception）、内存访问异常等。
	27	+
	28	+2. 内核代码错误：算子内核代码存在逻辑错误，如数组越界访问、空指针解引用、数据类型错误等。
	29	+
	30	+3. 内存问题：Device内存越界、地址对齐错误、未初始化内存访问等。
	31	+
	32	+4. 资源限制：任务超时、硬件资源不足、任务队列溢出等。
	33	+
	34	+## 处理步骤
	35	+
	36	+### 步骤1：定位失败的任务和内核
	37	+
	38	+从plog日志中提取关键信息，定位失败的任务：
	39	+
	40	+```bash
	41	+# 查找失败的任务信息
	42	+grep "Task run failed" plog.log
	43	+
	44	+# 输出示例：
	45	+# [ERROR] Task run failed, device_id=0, stream_id=2, task_id=1,
	46	+# sqe_type=0(ffts), errType=0x1(task exception), sqSwStatus=0
	47	+#
	48	+# 解读：
	49	+# - device_id=0: 设备ID为0
	50	+# - stream_id=2: 流ID为2
	51	+# - task_id=1: 任务ID为1
	52	+# - errType=0x1: 任务异常（task exception）
	53	+```
	54	+
	55	+```bash
	56	+# 查找失败的内核名称
	57	+grep "fault kernel_name" plog.log
	58	+
	59	+# 输出示例：
	60	+# fault kernel_name=Add_ee98c6628030785f610b924ab1557b31_high_performance_210000000
	61	+#
	62	+# 解读：
	63	+# - 内核名称为 Add_ee98c6628030785f610b924ab1557b31
	64	+# - 后缀 _high_performance 表示高性能模式
	65	+# - 210000000 为tiling key
	66	+```
	67	+
	68	+### 步骤2：分析错误码和错误类型
	69	+
	70	+从plog日志中提取错误码和错误类型：
	71	+
	72	+```bash
	73	+# 查找错误码
	74	+grep "retCode" plog.log
	75	+
	76	+```
	77	+
	78	+```bash
	79	+# 查找错误类型
	80	+grep "errType" plog.log
	81	+
	82	+# 常见错误类型：
	83	+# task exception - 任务异常
	84	+# task timeout - 任务超时
	85	+```
	86	+
	87	+### 步骤3：分析Device侧异常详情
	88	+
	89	+从plog日志中提取Device侧异常的详细信息：
	90	+
	91	+```bash
	92	+# 查找Device侧异常信息
	93	+grep "The error from device" plog.log
	94	+
	95	+# 输出示例：
	96	+# The error from device(chipId:3, dieId:0), serial number is 20,
	97	+# there is an aivec error exception, core id is 4, error code = 0,
	98	+#
	99	+# 解读：
	100	+# - chipId:3, dieId:0: 芯片和Die编号
	101	+# - aivec error exception: AI向量核异常
	102	+```
	103	+
	104	+### 步骤4：综合分析与问题定位
	105	+
	106	+综合以上信息进行问题定位：
	107	+
	108	+```bash
	109	+# 完整的错误定位流程
	110	+echo "=== Device侧异常定位分析 ==="
	111	+
	112	+echo -e "\n[1] 失败任务定位："
	113	+grep "Task run failed" plog.log \| tail -1
	114	+
	115	+echo -e "\n[2] 失败内核定位："
	116	+grep "fault kernel_name" plog.log \| tail -1
	117	+
	118	+echo -e "\n[3] 错误码分析："
	119	+grep "retCode\\|Kernel task happen error" plog.log \| tail -1
	120	+
	121	+echo -e "\n[4] Device异常详情："
	122	+grep "The error from device" plog.log \| tail -1
	123	+
	124	+echo -e "\n[5] 扩展错误信息："
	125	+grep "The extend info" plog.log \| tail -1
	126	+
	127	+echo -e "\n[6] 内核参数信息："
	128	+grep "AIC_INFO" plog.log \| tail -3
	129	+
	130	+echo -e "\n=== 分析完成 ==="
	131	+```
	132	+
	133	+典型问题诊断示例：
	134	+
	135	+场景1：地址访问越界
	136	+```
	137	+错误特征：
	138	+- aivec error exception
	139	+
	140	+诊断步骤：
	141	+1. 检查内核代码中的数组访问是否越界
	142	+2. 检查内存分配大小是否足够
	143	+3. 验证tiling参数计算是否正确
	144	+```
	145	+
	146	+场景2：内存对齐错误
	147	+```
	148	+错误特征：
	149	+- 数据传输失败
	150	+
	151	+诊断步骤：
	152	+1. 检查内存地址是否满足对齐要求（通常需要64字节对齐）
	153	+2. 检查数据结构定义是否正确
	154	+3. 验证内存分配接口的对齐参数
	155	+```
	156	+
	157	+场景3：任务超时
	158	+```
	159	+错误特征：
	160	+- task timeout
	161	+- 任务执行时间过长
	162	+
	163	+诊断步骤：
	164	+1. 检查内核是否存在死循环
	165	+2. 检查任务计算量是否过大
	166	+3. 验证硬件资源是否充足
	167	+```
	168	+
	169	+### 步骤5：启用详细日志
	170	+
	171	+如果默认日志信息不够，可以启用更详细的日志级别：
	172	+
	173	+```bash
	174	+# 设置日志级别（环境变量方式）
	175	+export ASCEND_GLOBAL_LOG_LEVEL=0 # DEBUG级别，输出最详细日志
	176	+export ASCEND_SLOG_PRINT_TO_STDOUT=1 # 输出到标准输出
	177	+
	178	+# 或在代码中设置
	179	+aclError error = aclInit(nullptr);
	180	+// Runtime会自动读取日志级别配置
	181	+
	182	+# 运行程序
	183	+./your_program 2>&1 \| tee detailed_plog.log
	184	+```
	185	+
	186	+日志级别说明：
	187	+- 0: DEBUG - 详细调试日志
	188	+- 1: INFO - 常规信息日志（默认）
	189	+- 2: WARNING - 错误和警告日志
	190	+- 3: ERROR - 仅错误日志
	191	+
	192	+## 相关 issue
	193	+
	194	+暂无相关Issue。

Ddocs/04_FAQ/析构函数中调用aclFinalize导致应用进程coredump.md+0-136

	@@ -1,136 +0,0 @@
1	-# 析构函数中调用去初始化接口aclFinalize导致应用进程coredump
2	-
3	-## 问题现象描述
4	-
5	-应用程序运行过程中出现core dump，应用程序异常终止。
6	-
7	-## 原因分析
8	-
9	-1. 生成coredump文件。
10	- - 物理机场景，执行ulimit -c unlimited命令，表示在程序崩溃时生成coredump文件：
11	-
12	- 完成问题定位后，如果不需要生成coredump文件，可执行ulimit -c 0命令。
13	-
14	- - Docker场景，在Docker启动命令中增加--ulimit core=-1设置。
15	-
16	-2. 运行应用程序，若进程崩溃，即可在当前目录下生成coredump文件。
17	-3. 使用gdb工具调试core文件、打印堆栈信息。
18	-
19	- 进入gdb模式，调试coredump文件，命令示例如下。其中，_main_表示产生coredump文件的可执行程序名称，可根据实际情况修改；coredump文件名需根据实际文件名称修改。
20	-
21	- ```
22	- gdb main core.
23	- ```
24	-
25	- 执行命令后，gdb工具会将发生异常的代码、其所在的函数、文件名和所在文件的行数打印到屏幕，堆栈信息的最上面是最底层的调用栈信息，方便定位问题。堆栈信息举例如下：
26	-
27	- ```
28	- Thread 1 "main" received signal SIGSEGV, Segmentation fault.
29	- 0x0000ffffa70747c8 in ge::PluginManager::~PluginManager() () from ******/lib64/libge_common.so
30	- (gdb) bt
31	- #0 0x0000ffffa70747c8 in ge::PluginManager::~PluginManager() () from ******/lib64/libge_common.so
32	- #1 0x0000ffffa707c900 in ge::RuntimePluginLoader::Finalize() () from ******/lib64/libge_common.so
33	- #2 0x0000ffffa29485d0 in ge::GeExecutor::FinalizeEx() () from ******/lib64/libge_executor.so
34	- #3 0x0000ffffb06fabc in aclFinalize() from ******/lib64/libascendcl.so
35	- #4 0x0000ffffbd5a98ec in ResourceManager::~ResourceManager() () from ******/envs/gly/lib/pythons3.7/site-packages/mindspore/_c_dataengine.cpython-37m-aarch64-linux-gnu.so
36	- #5 0x0000ffffbd5a9f80 in std::Sp_counted_ptr<ResourceManager, (__gnu_cxx::Lock_policy)2>::_M_dispose() () from *****/envs/gly/lib/pythons3.7/site-packages/mindspore/_c_dataengine.cpython-37m-aarch64-linux-gnu.so
37	- #6 0x0000ffffbd5a97f0 in std::shared_ptr<ResourceManager>::~shared_ptr() () from ******/envs/gly/lib/pythons3.7/site-packages/mindspore/_c_dataengine.cpython-37m-aarch64-linux-gnu.so
38	- ```
39	-
40	- 注意，调试coredump文件、打印堆栈信息要在出现问题的运行环境中，如果换一套环境，可能导致调试的堆栈信息不准确。
41	-
42	- 若环境中未安装gdb，则需要安装gdb，可通过包管理（如apt-get install gdb、yum install gdb）进行安装，详细安装步骤及使用方法请参见[GDB官方文档](https://sourceware.org/gdb/)。
43	-
44	-4. 分析堆栈信息。
45	-
46	- 生成coredump文件、检查打印的堆栈信息后，发现应用程序在调用aclFinalize接口后异常退出，因此初步判断可能是aclFinalize接口使用问题。
47	-
48	-5. 排查应用程序代码中aclFinalize接口的调用逻辑。
49	-
50	- 排查代码逻辑，发现该aclFinalize接口在析构函数中被调用，但该接口存在使用约束：不建议在析构函数中调用aclFinalize接口，否则在进程退出时可能由于单例析构顺序未知而导致进程异常退出的问题。因此判断是由于在析构函数中调用aclFinalize接口导致应用进程coredump。
51	-
52	-## 处理步骤
53	-
54	-优化应用程序的代码逻辑，不能在析构函数中调用aclFinalize接口，下文给出正确、错误的代码示例。
55	-
56	-- aclFinalize接口的正确调用示例如下：
57	-
58	- ```
59	- int main() {
60	- // 初始化
61	- // 此处的..表示相对路径，相对可执行文件所在的目录，例如，编译出来的可执行文件存放在out目录下，此处的..就表示out目录的上一级目录
62	- const char *aclConfigPath = "../src/acl.json";
63	- aclError ret = aclInit(aclConfigPath);
64	-
65	- // 业务处理代码
66	-
67	- // 去初始化，没有退出main函数，所有资源都可用
68	- ret = aclFinalize();
69	- return 0;
70	- }
71	- ```
72	-
73	-- aclFinalize接口的错误调用示例如下，使用单例析构去初始化：
74	-
75	- ```
76	- class ResourceManager {
77	- public:
78	- ResourceManager() = default;
79	- // 单例析构
80	- ~ResourceManager() {
81	- // 去初始化
82	- (void) aclFinalize();
83	- }
84	- // 单例构造
85	- static ResourceManager &Instance() {
86	- static ResourceManager instance;
87	- return instance;
88	- }
89	- aclError Init() {
90	- // 初始化
91	- // 此处的..表示相对路径，相对可执行文件所在的目录，例如，编译出来的可执行文件存放在out目录下，此处的..就表示out目录的上一级目录
92	- const char *aclConfigPath = "../src/acl.json";
93	- return aclInit(aclConfigPath);
94	- }
95	- };
96	- int main() {
97	- // 初始化
98	- aclError ret = ResourceManager::Instance().Init();
99	- // 业务处理代码
100	- // 没有显式去初始化，最后ResourceManager单例析构时调用aclFinalize
101	- // 由于单例析构是在main函数退出后才执行，单例析构和进程依赖so的卸载顺序无法控制
102	- // 会出现aclFinalize访问的一些资源所在so已经被卸载，从而导致进程退出异常
103	- return 0;
104	- }
105	- ```
106	-
107	-- aclFinalize接口的错误调用示例如下，使用全局变量析构去初始化：
108	-
109	- ```
110	- class ResourceManager {
111	- public:
112	- ResourceManager() = default;
113	- // 全局变量析构
114	- ~ResourceManager() {
115	- // 去初始化
116	- (void) aclFinalize();
117	- }
118	- aclError Init() {
119	- // 初始化
120	- // 此处的..表示相对路径，相对可执行文件所在的目录，例如，编译出来的可执行文件存放在out目录下，此处的..就表示out目录的上一级目录
121	- const char *aclConfigPath = "../src/acl.json";
122	- return aclInit(aclConfigPath);
123	- }
124	- };
125	- // 全局变量构造
126	- ResourceManager g_resource_manager;
127	- int main() {
128	- // 初始化
129	- aclError ret = g_resource_manager.Init();
130	- // 业务处理代码
131	- // 没有显式去初始化，最后ResourceManager全局变量析构时调用aclFinalize
132	- // 由于全局变量析构是在main函数退出后才执行，全局变量析构和进程依赖so的卸载顺序无法控制
133	- // 会出现aclFinalize访问的一些资源所在so已经被卸载，从而导致进程退出异常
134	- return 0;
135	- }
136	- ```

Ddocs/04_FAQ/用户进程异常退出后重启进程失败.md+0-37

	@@ -1,37 +0,0 @@
1	-# 用户进程异常退出后重启进程失败
2	-
3	-## 问题现象描述
4	-
5	-用户进程卡住或者用户强制退出进程后，再次重启，重启后发现进程无法正常启动。类似的日志信息如下：
6	-
7	-acl接口的报错信息：aclrtProcessReport failed
8	-
9	-```
10	-aclrtProcessReport failed, ret = 107012
11	-```
12	-
13	-Runtime日志信息：halResourceIdAlloc xxx failed
14	-
15	-```
16	-[ERROR] RUNTIME(2086,rtstest_host):2021-06-09-02:14:46.034.380 [npu_driver.cc:285]2086 StreamIdAlloc:[driver interface] halResourceIdAlloc streamid failed: device_id=0, tsId=0, drvRetCode=48!
17	-[ERROR] RUNTIME(2086,rtstest_host):2021-06-09-02:14:46.034.401 [stream.cc:448]2086 Setup:Failed to alloc stream id, retCode=0x702001a.
18	-[ERROR] RUNTIME(2086,rtstest_host):2021-06-09-02:14:46.034.416 [context.cc:1251]2086 StreamCreate:Setup stream failed, retCode=0x702001a.
19	-[ERROR] RUNTIME(2086,rtstest_host):2021-06-09-02:14:46.034.440 [logger.cc:211]2086 StreamCreate:Create stream failed, priority=7 ,flags=0.
20	-[ERROR] RUNTIME(2086,rtstest_host):2021-06-09-02:14:46.034.458 [api_c.cc:461]2086 rtStreamCreateWithFlags:ErrCode=207008, desc=[driver error:no stream resource], InnerCode=0x702001a
21	-[ERROR] RUNTIME(2086,rtstest_host):2021-06-09-02:14:46.034.469 [error_message_manage.cc:26]2086 ReportFuncErrorReason:rtStreamCreateWithFlags execute failed, reason=[driver error:no stream resource]
22	-```
23	-
24	-## 可能原因
25	-
26	-通过日志分析无法正常重启的原因可能是public taskid、stream id、eventid等资源申请不到引起的：
27	-
28	-- 资源已经被其他进程占用完。
29	-- 上一个进程退出时还未完全释放完资源。
30	-
31	-## 处理步骤
32	-
33	-针对上述可能原因，可以按以下方式处理：
34	-
35	-- 等待一分钟后再重新启动进程，保证上一个进程资源释放完成。
36	-- 停止其他进程或者等其他进程执行完成后再启动进程。
37	-- 如果通过上述方式处理后仍然申请失败，建议检查是否超过了可用的资源上限，如果未超上限，则需要重启环境强行释放资源、恢复环境。

Adocs/04_FAQ/算子执行输出全0的常见原因排查.md+123-0

@@ -0,0 +1,123 @@
	1	+# 算子执行输出全0的常见原因排查
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：执行算子后输出数据全为0
	6	+
	7	+执行算子时，算子输入数据正确，但输出数据异常，全为0。
	8	+
	9	+报错日志示例如下：
	10	+```
	11	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.036.721 [stars_engine.cc:1321]2082291 ProcLogicCqReport:[INIT][DEFAULT]Task run failed, device_id=0, stream_id=2, task_id=1, sqe_type=0(ffts), errType=0x1(task exception), sqSwStatus=0
	12	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.365 [davinic_kernel_task.cc:1219]2082291 PreCheckTaskErr:[INIT][DEFAULT]Kernel task happen error, retCode=0x31, [vector core exception].
	13	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.050.517 [stream.cc:1082]2082291 GetError:[INIT][DEFAULT]AIV Kernel happen error, retCode=0x31.
	14	+[ERROR] RUNTIME(2082291,python3):2024-07-04-14:14:25.051.013 [davinic_kernel_task.cc:1082]2082291 GetArgsInfo:[INIT][DEFAULT][AIC_INFO] args(0 to 9) after execute:0, 0, 0, 0, 0, 0, 0, 0, 0
	15	+```
	16	+
	17	+## 可能原因
	18	+
	19	+1. 内存未初始化：输出内存未正确初始化或清零，导致读取到未初始化的数据（通常为0）。
	20	+2. 数据类型不匹配：算子输入输出数据类型配置错误，例如float32与float16混用，导致数据解析错误。
	21	+3. 内存拷贝失败：Host与Device之间的内存拷贝操作失败或未执行，导致Device侧内存为初始值（0）。
	22	+4. 内核编译或执行异常：
	23	+ - 低版本内核（5.10以下）使用asan工具导致执行异常。
	24	+ - 算子编译选项不正确。
	25	+ - 内核启动参数配置错误。
	26	+
	27	+5. 地址访问越界：算子访问了非法内存地址，触发Device侧异常。
	28	+
	29	+## 处理步骤
	30	+
	31	+### 排查步骤1：检查内存初始化
	32	+
	33	+确保输出内存已正确分配和初始化：
	34	+
	35	+```cpp
	36	+// 检查内存分配是否成功
	37	+void *devPtr = nullptr;
	38	+size_t devSize = 1024 * 1024; // 1MB
	39	+aclError error = aclrtMalloc(&devPtr, devSize, ACL_MEM_MALLOC_HUGE_FIRST);
	40	+if (error != ACL_RT_SUCCESS \|\| devPtr == nullptr) {
	41	+ printf("Memory allocation failed\n");
	42	+ return error;
	43	+}
	44	+
	45	+// 初始化内存为特定值（非0）用于验证
	46	+error = aclrtMemset(devPtr, devSize, 0xFF, devSize);
	47	+if (error != ACL_RT_SUCCESS) {
	48	+ printf("Memory initialization failed\n");
	49	+ aclrtFree(devPtr);
	50	+ return error;
	51	+}
	52	+
	53	+// 执行算子...
	54	+// 执行后检查输出，如果仍为0xFF说明算子未写入
	55	+// 如果为0说明算子写入但数据异常
	56	+```
	57	+
	58	+### 排查步骤2：验证数据类型匹配
	59	+
	60	+确保算子输入输出数据类型一致：
	61	+
	62	+```cpp
	63	+// 检查tensor描述符的数据类型
	64	+aclDataType inputType, outputType;
	65	+aclGetTensorDescType(inputDesc, &inputType);
	66	+aclGetTensorDescType(outputDesc, &outputType);
	67	+
	68	+if (inputType != outputType) {
	69	+ printf("Warning: Input type %d != Output type %d\n", inputType, outputType);
	70	+}
	71	+
	72	+// 确认数据格式（ND、NCHW等）匹配
	73	+aclFormat inputFormat, outputFormat;
	74	+aclGetTensorDescFormat(inputDesc, &inputFormat);
	75	+aclGetTensorDescFormat(outputDesc, &outputFormat);
	76	+
	77	+printf("Input format: %d, Output format: %d\n", inputFormat, outputFormat);
	78	+```
	79	+
	80	+### 排查步骤3：检查内存拷贝操作
	81	+
	82	+验证Host与Device之间的数据传输是否成功：
	83	+
	84	+```cpp
	85	+// Host -> Device拷贝
	86	+aclError error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	87	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	88	+if (error != ACL_RT_SUCCESS) {
	89	+ printf("Host to Device copy failed: %d\n", error);
	90	+ return error;
	91	+}
	92	+
	93	+// 确保拷贝完成
	94	+error = aclrtSynchronizeStream(stream);
	95	+if (error != ACL_RT_SUCCESS) {
	96	+ printf("Synchronize failed after H2D copy: %d\n", error);
	97	+ return error;
	98	+}
	99	+
	100	+// 可选：拷贝回Host验证数据是否正确传输
	101	+error = aclrtMemcpyAsync(hostVerifyPtr, hostSize, devPtr, devSize,
	102	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	103	+error = aclrtSynchronizeStream(stream);
	104	+
	105	+// 比较hostPtr和hostVerifyPtr的数据是否一致
	106	+```
	107	+
	108	+### 排查步骤4：分析plog日志定位异常
	109	+
	110	+通过plog日志中的错误码和异常信息定位问题：
	111	+```bash
	112	+# 查找关键错误信息
	113	+grep -i "error\\|exception\\|failed" plog.log
	114	+
	115	+# 定位错误内核
	116	+grep "fault kernel_name" plog.log
	117	+# 示例输出：fault kernel_name=Add_ee98c6628030785f610b924ab1557b31
	118	+
	119	+```
	120	+
	121	+## 相关 issue
	122	+
	123	+- [Issue #480: aicore异常报错导致算子执行失败](https://gitcode.com/cann/runtime/issues/480)

Adocs/04_FAQ/跨DeviceP2P数据交互配置失败.md+191-0

@@ -0,0 +1,191 @@
	1	+# 跨Device P2P数据交互配置失败
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：跨Device内存复制失败
	6	+
	7	+在多Device场景下，尝试进行跨Device的P2P数据交互时，内存复制失败。
	8	+
	9	+错误代码示例：
	10	+```cpp
	11	+aclInit(NULL);
	12	+aclrtSetDevice(0);
	13	+void *dev0Mem = nullptr;
	14	+aclrtMalloc(&dev0Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST);
	15	+
	16	+aclrtSetDevice(1);
	17	+void *dev1Mem = nullptr;
	18	+aclrtMalloc(&dev1Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST);
	19	+
	20	+// 尝试跨Device内存复制失败
	21	+aclrtMemcpy(dev1Mem, 1024, dev0Mem, 1024, ACL_MEMCPY_DEVICE_TO_DEVICE); // ✗ 失败
	22	+```
	23	+
	24	+现象2：调用aclrtDeviceEnablePeerAccess接口报错
	25	+
	26	+调用`aclrtDeviceEnablePeerAccess`接口时报错，提示Device间不支持P2P访问。
	27	+
	28	+## 可能原因
	29	+
	30	+1. 未检查Device间的P2P访问能力：并非所有Device组合都支持P2P数据交互，需要先使用`aclrtDeviceCanAccessPeer`接口查询。
	31	+2. 未正确开启P2P访问权限：即使Device间支持P2P，也需要显式调用`aclrtDeviceEnablePeerAccess`接口开启访问权限。
	32	+3. 内存分配未使用P2P优化标志：跨Device内存复制时，应使用`ACL_MEM_MALLOC_HUGE_FIRST_P2P`标志分配内存以获得更好的性能。
	33	+4. 硬件组网限制：Device之间需要处于PCIe等互联的组网拓扑下才能支持P2P访问。
	34	+
	35	+## 处理步骤
	36	+
	37	+方案一：检查Device间P2P访问能力并开启访问权限
	38	+
	39	+按照正确的流程配置P2P数据交互：
	40	+
	41	+```cpp
	42	+aclInit(NULL);
	43	+
	44	+// 1. 检查Device 0和Device 1之间是否支持P2P
	45	+int32_t canAccessPeer = 0;
	46	+aclrtDeviceCanAccessPeer(&canAccessPeer, 0, 1); // 查询Device 0能否访问Device 1
	47	+
	48	+if (canAccessPeer == 1) {
	49	+ // 2. 开启Device 0到Device 1的P2P访问
	50	+ aclrtSetDevice(0);
	51	+ uint32_t reserveFlag = 0U;
	52	+ aclrtDeviceEnablePeerAccess(1, reserveFlag); // 开启Device 0→Device 1的访问
	53	+
	54	+ void *dev0Mem = nullptr;
	55	+ aclrtMalloc(&dev0Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
	56	+
	57	+ // 3. 开启Device 1到Device 0的P2P访问
	58	+ aclrtSetDevice(1);
	59	+ aclrtDeviceEnablePeerAccess(0, reserveFlag); // 开启Device 1→Device 0的访问
	60	+
	61	+ void *dev1Mem = nullptr;
	62	+ aclrtMalloc(&dev1Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
	63	+
	64	+ // 4. 执行跨Device内存复制
	65	+ aclrtMemcpy(dev1Mem, 1024, dev0Mem, 1024, ACL_MEMCPY_DEVICE_TO_DEVICE); // ✓ 成功
	66	+
	67	+ // 5. 关闭P2P访问并释放资源
	68	+ aclrtDeviceDisablePeerAccess(0); // 关闭Device 1→Device 0的访问
	69	+ aclrtFree(dev1Mem);
	70	+ aclrtResetDevice(1);
	71	+
	72	+ aclrtSetDevice(0);
	73	+ aclrtDeviceDisablePeerAccess(1); // 关闭Device 0→Device 1的访问
	74	+ aclrtFree(dev0Mem);
	75	+ aclrtResetDeviceForce(0);
	76	+} else {
	77	+ printf("Device 0 and Device 1 don't support P2P feature\n");
	78	+}
	79	+
	80	+aclFinalize();
	81	+```
	82	+
	83	+方案二：使用Host内存作为中转
	84	+
	85	+如果Device间不支持P2P或P2P配置失败，可以使用Host内存作为中转：
	86	+
	87	+```cpp
	88	+aclrtSetDevice(0);
	89	+void *dev0Mem = nullptr;
	90	+aclrtMalloc(&dev0Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST);
	91	+
	92	+// 分配Host内存作为中转
	93	+void *hostMem = nullptr;
	94	+aclrtMallocHost(&hostMem, 1024);
	95	+
	96	+// Device 0 → Host
	97	+aclrtMemcpy(hostMem, 1024, dev0Mem, 1024, ACL_MEMCPY_DEVICE_TO_HOST);
	98	+
	99	+aclrtSetDevice(1);
	100	+void *dev1Mem = nullptr;
	101	+aclrtMalloc(&dev1Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST);
	102	+
	103	+// Host → Device 1
	104	+aclrtMemcpy(dev1Mem, 1024, hostMem, 1024, ACL_MEMCPY_HOST_TO_DEVICE);
	105	+
	106	+aclrtFreeHost(hostMem);
	107	+aclrtFree(dev1Mem);
	108	+aclrtResetDevice(1);
	109	+
	110	+aclrtSetDevice(0);
	111	+aclrtFree(dev0Mem);
	112	+aclrtResetDeviceForce(0);
	113	+```
	114	+
	115	+方案三：检查硬件组网拓扑
	116	+
	117	+确认Device间的硬件连接状态：
	118	+
	119	+```cpp
	120	+uint32_t deviceCount;
	121	+aclrtGetDeviceCount(&deviceCount);
	122	+
	123	+// 检查所有Device组合的P2P能力
	124	+for (uint32_t i = 0; i < deviceCount; ++i) {
	125	+ for (uint32_t j = 0; j < deviceCount; ++j) {
	126	+ if (i != j) {
	127	+ int32_t canAccessPeer = 0;
	128	+ aclrtDeviceCanAccessPeer(&canAccessPeer, i, j);
	129	+ printf("Device %d can access Device %d: %s\n",
	130	+ i, j, canAccessPeer ? "Yes" : "No");
	131	+ }
	132	+ }
	133	+}
	134	+```
	135	+
	136	+方案四：完整示例代码
	137	+
	138	+完整的P2P数据交互流程示例：
	139	+
	140	+```cpp
	141	+#include <stdio.h>
	142	+#include "acl/acl.h"
	143	+
	144	+int main() {
	145	+ aclInit(NULL);
	146	+
	147	+ int32_t canAccessPeer = 0;
	148	+ aclrtDeviceCanAccessPeer(&canAccessPeer, 0, 1);
	149	+
	150	+ if (canAccessPeer == 1) {
	151	+ // Device 0操作
	152	+ aclrtSetDevice(0);
	153	+ uint32_t reserveFlag = 0U;
	154	+ aclrtDeviceEnablePeerAccess(1, reserveFlag);
	155	+
	156	+ void *dev0Mem = nullptr;
	157	+ aclrtMalloc(&dev0Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
	158	+
	159	+ // Device 1操作
	160	+ aclrtSetDevice(1);
	161	+ aclrtDeviceEnablePeerAccess(0, reserveFlag);
	162	+
	163	+ void *dev1Mem = nullptr;
	164	+ aclrtMalloc(&dev1Mem, 1024, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
	165	+
	166	+ // 跨Device内存复制
	167	+ aclrtMemcpy(dev1Mem, 1024, dev0Mem, 1024, ACL_MEMCPY_DEVICE_TO_DEVICE);
	168	+
	169	+ // 清理资源
	170	+ aclrtDeviceDisablePeerAccess(0);
	171	+ aclrtFree(dev1Mem);
	172	+ aclrtResetDevice(1);
	173	+
	174	+ aclrtSetDevice(0);
	175	+ aclrtDeviceDisablePeerAccess(1);
	176	+ aclrtFree(dev0Mem);
	177	+ aclrtResetDeviceForce(0);
	178	+
	179	+ printf("P2P copy success\n");
	180	+ } else {
	181	+ printf("Current device doesn't support P2P feature\n");
	182	+ }
	183	+
	184	+ aclFinalize();
	185	+ return 0;
	186	+}
	187	+```
	188	+
	189	+## 相关 issue
	190	+
	191	+暂无相关Issue。

Adocs/04_FAQ/进程间IPC内存共享的页表对齐要求.md+121-0

@@ -0,0 +1,121 @@
	1	+# 进程间IPC内存共享的页表对齐要求
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：调用IPC接口报错提示页表未对齐
	6	+
	7	+在使用Runtime IPC接口进行进程间内存共享时，调用`aclrtIpcMemGetExportKey`或`aclrtIpcMemImportByKey`接口时报错，提示共享内存页表未对齐。
	8	+
	9	+错误代码示例：
	10	+```cpp
	11	+// 进程A：分配内存并导出IPC key
	12	+void *ptrA = nullptr;
	13	+aclrtSetDevice(0);
	14	+aclrtMalloc(&ptrA, 1024); // 分配1024字节内存
	15	+
	16	+char key[65];
	17	+aclrtIpcMemGetExportKey(ptrA, 1024, key, 65, ACL_RT_IPC_MEM_EXPORT_FLAG_DISABLE_PID_VALIDATION);
	18	+// ✗ 可能失败：内存未页表对齐
	19	+```
	20	+
	21	+## 可能原因
	22	+
	23	+1. 内存分配未对齐：通过`aclrtMalloc`接口分配设备内存时，出于性能考虑，可能会从更大的底层内存块中切分出来，导致分配的内存地址未页表对齐。
	24	+2. 页表对齐安全检查：IPC接口会检查共享内存是否页表对齐，若未对齐，API将拦截并报错，以防止跨进程多映射内存导致的信息泄露风险。
	25	+3. 不同内存类型页表大小不同：普通页内存的页表大小为4K，大页内存的页表大小支持2M或1G。
	26	+
	27	+## 处理步骤
	28	+
	29	+方案一：使用aclrtMalloc接口的内存分配规则
	30	+
	31	+按照内存分配规则申请内存，确保内存页表对齐：
	32	+
	33	+```cpp
	34	+// 正确示例：使用合适的内存分配标志
	35	+void *ptrA = nullptr;
	36	+aclrtSetDevice(0);
	37	+
	38	+// 对于小内存（小于2M），使用普通页内存分配
	39	+aclrtMalloc(&ptrA, 4096, ACL_MEM_MALLOC_HUGE_FIRST); // ✓ 分配4K对齐的内存
	40	+
	41	+char key[65];
	42	+aclrtIpcMemGetExportKey(ptrA, 4096, key, 65, ACL_RT_IPC_MEM_EXPORT_FLAG_DISABLE_PID_VALIDATION); // ✓ 成功
	43	+```
	44	+
	45	+方案二：分配大页内存以确保对齐
	46	+
	47	+对于需要大页内存的场景，使用大页内存分配标志：
	48	+
	49	+```cpp
	50	+// 正确示例：分配大页内存（2M或1G对齐）
	51	+void *ptrA = nullptr;
	52	+aclrtSetDevice(0);
	53	+
	54	+// 分配2M大页内存
	55	+size_t largePageSize = 2 * 1024 * 1024; // 2MB
	56	+aclrtMalloc(&ptrA, largePageSize, ACL_MEM_MALLOC_HUGE_ONLY); // ✓ 强制使用大页内存
	57	+
	58	+char key[65];
	59	+aclrtIpcMemGetExportKey(ptrA, largePageSize, key, 65, ACL_RT_IPC_MEM_EXPORT_FLAG_DISABLE_PID_VALIDATION);
	60	+```
	61	+
	62	+方案三：使用VMM接口实现进程间共享
	63	+
	64	+如果无法保证内存对齐，可以使用VMM（Virtual Memory Management）接口，它提供更灵活的内存管理和共享功能：
	65	+
	66	+```cpp
	67	+// 进程A：使用VMM接口
	68	+const size_t dataSize = 1024 * sizeof(float);
	69	+aclrtPhysicalMemProp prop = {};
	70	+prop.handleType = ACL_MEM_HANDLE_TYPE_NONE;
	71	+prop.allocationType = ACL_MEM_ALLOCATION_TYPE_PINNED;
	72	+prop.location.type = ACL_MEM_LOCATION_TYPE_DEVICE;
	73	+prop.location.id = 0;
	74	+prop.memAttr = ACL_HBM_MEM_NORMAL;
	75	+
	76	+// 查询内存申请粒度
	77	+size_t granularity = 0UL;
	78	+aclrtMemGetAllocationGranularity(&prop, ACL_RT_MEM_ALLOC_GRANULARITY_MINIMUM, &granularity);
	79	+
	80	+// 基于内存申请粒度申请物理内存（自动对齐）
	81	+size_t alignedSize = ((dataSize + granularity - 1U) / granularity) * granularity;
	82	+aclrtDrvMemHandle handle = nullptr;
	83	+aclrtMallocPhysical(&handle, alignedSize, &prop, 0);
	84	+
	85	+// 预留虚拟内存
	86	+void *virPtr;
	87	+aclrtReserveMemAddress(&virPtr, alignedSize, 0, nullptr, 0);
	88	+
	89	+// 将虚拟内存映射到物理内存
	90	+aclrtMapMem(virPtr, alignedSize, 0, handle, 0);
	91	+
	92	+// 导出共享handle
	93	+uint64_t shareableHandle = 0ULL;
	94	+aclrtMemExportToShareableHandle(handle, ACL_MEM_HANDLE_TYPE_NONE,
	95	+ ACL_RT_VMM_EXPORT_FLAG_DISABLE_PID_VALIDATION, &shareableHandle);
	96	+```
	97	+
	98	+方案四：检查内存对齐状态
	99	+
	100	+在导出IPC key前，可以先检查内存对齐状态：
	101	+
	102	+```cpp
	103	+void *ptrA = nullptr;
	104	+aclrtSetDevice(0);
	105	+aclrtMalloc(&ptrA, size, ACL_MEM_MALLOC_HUGE_FIRST);
	106	+
	107	+// 检查内存地址是否对齐（示例：4K对齐）
	108	+if ((uintptr_t)ptrA % 4096 != 0) {
	109	+ // 内存未对齐，需要调整分配策略或使用VMM接口
	110	+ printf("Warning: Memory not page-aligned, consider using VMM interface\n");
	111	+ aclrtFree(ptrA);
	112	+ // 切换到VMM方案或调整分配大小
	113	+} else {
	114	+ char key[65];
	115	+ aclrtIpcMemGetExportKey(ptrA, size, key, 65, ACL_RT_IPC_MEM_EXPORT_FLAG_DISABLE_PID_VALIDATION);
	116	+}
	117	+```
	118	+
	119	+## 相关 issue
	120	+
	121	+- [Issue #148: Stream有序内存分配IPC内存池共享](https://gitcode.com/cann/runtime/issues/148)

Adocs/04_FAQ/遇错即停模式下错误定位方法.md+305-0

@@ -0,0 +1,305 @@
	1	+# 遇错即停模式下错误定位方法
	2	+
	3	+## 问题现象描述
	4	+
	5	+现象1：遇错即停模式下无法准确定位失败任务
	6	+
	7	+在Stream上配置了遇错即停模式（ACL_STOP_ON_FAILURE），任务执行失败时整个Context下的任务被停止，但无法准确定位是哪个任务失败，失败原因是什么。
	8	+
	9	+典型场景：
	10	+```cpp
	11	+aclrtStream stream;
	12	+aclrtCreateStream(&stream);
	13	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	14	+
	15	+// 下发多个任务
	16	+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	17	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	18	+kernel1<<<8, nullptr, stream>>>(args1);
	19	+kernel2<<<8, nullptr, stream>>>(args2);
	20	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	21	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	22	+
	23	+// 同步时发现错误，但不知道是哪个任务失败
	24	+aclError error = aclrtSynchronizeStream(stream);
	25	+if (error != ACL_RT_SUCCESS) {
	26	+ // 如何定位是kernel1还是kernel2失败？
	27	+ // 失败原因是什么？
	28	+}
	29	+```
	30	+
	31	+## 可能原因
	32	+
	33	+1. 任务序列较长：Stream上下发了多个异步任务，错误发生后无法直观判断是哪个任务失败。
	34	+
	35	+2. 错误信息不够详细：仅通过返回码无法获知具体失败原因，需要结合plog日志和错误查询接口。
	36	+
	37	+3. 遇错即停特性限制：遇错即停模式会在首个错误发生时停止Context中所有Stream的任务执行，可能影响错误现场的保护。
	38	+
	39	+## 处理步骤
	40	+
	41	+### 方法1：分段同步定位错误任务
	42	+
	43	+将长任务序列分解为多个小段，每段后进行同步检查：
	44	+
	45	+```cpp
	46	+aclrtStream stream;
	47	+aclrtCreateStream(&stream);
	48	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	49	+
	50	+// 任务1：Host -> Device拷贝
	51	+aclError error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	52	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	53	+if (error != ACL_RT_SUCCESS) {
	54	+ printf("Task 1 launch failed: %d\n", error);
	55	+ return error;
	56	+}
	57	+
	58	+error = aclrtSynchronizeStream(stream);
	59	+if (error != ACL_RT_SUCCESS) {
	60	+ printf("Task 1 execution failed: %d\n", error);
	61	+ char *errMsg = aclGetRecentErrMsg();
	62	+ printf("Error: %s\n", errMsg);
	63	+ return error;
	64	+}
	65	+
	66	+// 任务2：kernel1执行
	67	+kernel1<<<8, nullptr, stream>>>(args1);
	68	+error = aclrtSynchronizeStream(stream);
	69	+if (error != ACL_RT_SUCCESS) {
	70	+ printf("Task 2 (kernel1) execution failed: %d\n", error);
	71	+ char *errMsg = aclGetRecentErrMsg();
	72	+ printf("Error: %s\n", errMsg);
	73	+ return error;
	74	+}
	75	+
	76	+// 任务3：kernel2执行
	77	+kernel2<<<8, nullptr, stream>>>(args2);
	78	+error = aclrtSynchronizeStream(stream);
	79	+if (error != ACL_RT_SUCCESS) {
	80	+ printf("Task 3 (kernel2) execution failed: %d\n", error);
	81	+ char *errMsg = aclGetRecentErrMsg();
	82	+ printf("Error: %s\n", errMsg);
	83	+ return error;
	84	+}
	85	+
	86	+// 任务4：Device -> Host拷贝
	87	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	88	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	89	+error = aclrtSynchronizeStream(stream);
	90	+if (error != ACL_RT_SUCCESS) {
	91	+ printf("Task 4 execution failed: %d\n", error);
	92	+ return error;
	93	+}
	94	+
	95	+aclrtDestroyStream(stream);
	96	+```
	97	+
	98	+### 方法2：使用Event标记任务边界
	99	+
	100	+在关键任务之间插入Event，通过Event状态判断任务执行进度：
	101	+
	102	+```cpp
	103	+aclrtStream stream;
	104	+aclrtCreateStream(&stream);
	105	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	106	+
	107	+// 创建多个Event用于标记任务边界
	108	+aclrtEvent event1, event2, event3;
	109	+aclrtCreateEvent(&event1);
	110	+aclrtCreateEvent(&event2);
	111	+aclrtCreateEvent(&event3);
	112	+
	113	+// 下发任务并插入Event标记
	114	+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	115	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	116	+aclrtRecordEvent(event1, stream); // 标记拷贝任务完成
	117	+
	118	+kernel1<<<8, nullptr, stream>>>(args1);
	119	+aclrtRecordEvent(event2, stream); // 标记kernel1完成
	120	+
	121	+kernel2<<<8, nullptr, stream>>>(args2);
	122	+aclrtRecordEvent(event3, stream); // 标记kernel2完成
	123	+
	124	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	125	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	126	+
	127	+// 同步并检查错误
	128	+aclError error = aclrtSynchronizeStream(stream);
	129	+if (error != ACL_RT_SUCCESS) {
	130	+ // 查询Event状态，定位失败任务
	131	+ aclrtEventStatus status1, status2, status3;
	132	+
	133	+ aclrtEventQuery(event1, &status1);
	134	+ aclrtEventQuery(event2, &status2);
	135	+ aclrtEventQuery(event3, &status3);
	136	+
	137	+ printf("Event status: event1=%d, event2=%d, event3=%d\n",
	138	+ status1, status2, status3);
	139	+
	140	+ // ACL_RT_EVENT_COMPLETE = 1, ACL_RT_EVENT_NOT_READY = 0
	141	+ // 最后一个完成的Event之后的任务为失败任务
	142	+ if (status1 == ACL_RT_EVENT_COMPLETE && status2 == ACL_RT_EVENT_NOT_READY) {
	143	+ printf("Task failed between event1 and event2 (kernel1 failed)\n");
	144	+ } else if (status2 == ACL_RT_EVENT_COMPLETE && status3 == ACL_RT_EVENT_NOT_READY) {
	145	+ printf("Task failed between event2 and event3 (kernel2 failed)\n");
	146	+ }
	147	+
	148	+ char *errMsg = aclGetRecentErrMsg();
	149	+ printf("Error detail: %s\n", errMsg);
	150	+}
	151	+
	152	+// 清理资源
	153	+aclrtDestroyEvent(event1);
	154	+aclrtDestroyEvent(event2);
	155	+aclrtDestroyEvent(event3);
	156	+aclrtDestroyStream(stream);
	157	+```
	158	+
	159	+### 方法3：结合错误查询接口和plog日志
	160	+
	161	+通过错误查询接口获取错误码，结合plog日志定位具体失败原因：
	162	+
	163	+```cpp
	164	+aclrtStream stream;
	165	+aclrtCreateStream(&stream);
	166	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	167	+
	168	+// 下发任务
	169	+// ...（省略任务下发代码）
	170	+
	171	+// 同步并检查错误
	172	+aclError error = aclrtSynchronizeStream(stream);
	173	+if (error != ACL_RT_SUCCESS) {
	174	+ // 1. 获取错误码
	175	+ printf("Stream synchronize failed with error code: %d (0x%x)\n", error, error);
	176	+
	177	+ // 2. 查看最近错误
	178	+ aclError lastError = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL);
	179	+ printf("Last error: %d (0x%x)\n", lastError, lastError);
	180	+
	181	+ // 3. 获取详细错误信息
	182	+ char *errMsg = aclGetRecentErrMsg();
	183	+ if (errMsg != nullptr) {
	184	+ printf("Error message: %s\n", errMsg);
	185	+ }
	186	+
	187	+ // 4. 提示查看plog日志
	188	+ printf("Please check plog for detailed error information:\n");
	189	+ printf(" - Look for 'fault kernel_name' to identify failed kernel\n");
	190	+ printf(" - Look for 'retCode' to get error code\n");
	191	+ printf(" - Look for 'error module_type' and 'module_name' for error source\n");
	192	+}
	193	+
	194	+aclrtDestroyStream(stream);
	195	+```
	196	+
	197	+plog日志关键信息提取：
	198	+```bash
	199	+# 查找失败的任务
	200	+grep "Task run failed" plog.log
	201	+
	202	+# 查找失败的内核名称
	203	+grep "fault kernel_name" plog.log
	204	+
	205	+# 查找错误码和类型
	206	+grep "retCode\\|errType" plog.log
	207	+
	208	+# 查找错误模块
	209	+grep "module_type\\|module_name" plog.log
	210	+
	211	+# 示例输出解析：
	212	+# [ERROR] Task run failed, device_id=0, stream_id=2, task_id=1
	213	+# -> 定位到具体的设备、流、任务ID
	214	+# fault kernel_name=Add_ee98c6628030785f610b924ab1557b31
	215	+# -> 定位到具体失败的内核
	216	+# retCode=0x31, [vector core exception]
	217	+# -> 错误原因：向量核异常
	218	+# module_type=5, module_name=EZ9999
	219	+# -> 错误来源模块
	220	+```
	221	+
	222	+### 方法4：对比遇错继续模式
	223	+
	224	+如果遇错即停模式难以定位，可临时切换到遇错继续模式，让后续任务继续执行以收集更多信息：
	225	+
	226	+```cpp
	227	+aclrtStream stream;
	228	+aclrtCreateStream(&stream);
	229	+
	230	+// 测试时使用遇错继续模式
	231	+aclrtSetStreamFailureMode(stream, ACL_CONTINUE_ON_FAILURE);
	232	+
	233	+// 下发多个任务
	234	+// ...
	235	+
	236	+// 同步并检查错误
	237	+aclError error = aclrtSynchronizeStream(stream);
	238	+if (error != ACL_RT_SUCCESS) {
	239	+ printf("Error occurred: %d\n", error);
	240	+
	241	+ // 遇错继续模式下，可以获取更多信息
	242	+ // 但注意：aclrtPeekAtLastError可能返回的是最后发生的错误，而非首次错误
	243	+ aclError lastError = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL);
	244	+ printf("Last error: %d\n", lastError);
	245	+
	246	+ char *errMsg = aclGetRecentErrMsg();
	247	+ printf("Error message: %s\n", errMsg);
	248	+
	249	+ // 查看哪些任务执行成功，哪些失败
	250	+ // （需要通过中间结果验证）
	251	+}
	252	+
	253	+// 问题定位后，改回遇错即停模式用于生产环境
	254	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	255	+aclrtDestroyStream(stream);
	256	+```
	257	+
	258	+### 方法5：使用Host回调任务进行调试
	259	+
	260	+在关键任务之间插入Host回调函数，记录任务执行进度：
	261	+
	262	+```cpp
	263	+void taskCallback(void *userData) {
	264	+ int taskId = (int)userData;
	265	+ printf("Task %d completed successfully\n", *taskId);
	266	+}
	267	+
	268	+// ...
	269	+
	270	+aclrtStream stream;
	271	+aclrtCreateStream(&stream);
	272	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	273	+
	274	+int taskId1 = 1, taskId2 = 2, taskId3 = 3, taskId4 = 4;
	275	+
	276	+// 下发任务并插入回调
	277	+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize,
	278	+ ACL_MEMCPY_HOST_TO_DEVICE, stream);
	279	+aclrtLaunchHostFunc(stream, taskCallback, &taskId1);
	280	+
	281	+kernel1<<<8, nullptr, stream>>>(args1);
	282	+aclrtLaunchHostFunc(stream, taskCallback, &taskId2);
	283	+
	284	+kernel2<<<8, nullptr, stream>>>(args2);
	285	+aclrtLaunchHostFunc(stream, taskCallback, &taskId3);
	286	+
	287	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize,
	288	+ ACL_MEMCPY_DEVICE_TO_HOST, stream);
	289	+aclrtLaunchHostFunc(stream, taskCallback, &taskId4);
	290	+
	291	+// 同步
	292	+aclError error = aclrtSynchronizeStream(stream);
	293	+if (error != ACL_RT_SUCCESS) {
	294	+ // 通过回调输出可以知道最后一个成功完成的任务
	295	+ // 失败的任务就是最后一个成功任务之后的任务
	296	+ printf("Error occurred after last successful task\n");
	297	+}
	298	+
	299	+aclrtDestroyStream(stream);
	300	+```
	301	+
	302	+## 相关 issue
	303	+
	304	+- [Issue #487: ACL Graph捕获阶段错误处理咨询](https://gitcode.com/cann/runtime/issues/487)
	305	+- [Issue #544: 错误码格式及日志宏使用规范](https://gitcode.com/cann/runtime/issues/544)

Mdocs/README.md+29-175

	@@ -12,47 +12,13 @@
12	<!-- ==================== 文档导航 ==================== -->	12	<!-- ==================== 文档导航 ==================== -->
13	## 📚 文档导航	13	## 📚 文档导航
14		14
15	-<table width="100%" style="table-layout: fixed; width: 100%; border-collapse: collapse;">	15	+\| 文档 \| 定位与内容 \| 入口 \|
16	-<colgroup>	16	+\|------\|-----------\|------\|
17	-<col width="15%">	17	+\| ⚡ 快速入门 \| 零基础起步，第一次接触 Runtime 的开发者。Runtime 简介 + 编程模型讲解，建议顺序阅读 \| [Runtime 简介](01_quick_start/Runtime简介.md) · [编程模型](01_quick_start/Runtime编程模型.md) \|
18	-<col width="55%">	18	+\| 📖 编程指南 \| 深度开发手册，已跑通入门阶段 Hello CANN 后需要深入理解原理的开发者。原理解析 + 示例代码 \| [编程指南](02_dev_guide/00_dev_guide.md) \|
19	-<col width="30%">	19	+\| 📋 API 参考 \| 接口字典，开发过程中随时查阅函数定义。函数签名 + 参数说明 + 返回值 \| [头文件说明](03_api_ref/01_概述.md) · [API 参考](03_api_ref/api_ref.md) \|
20	-</colgroup>	20	+\| 🏗️ 架构指南 \| 面向贡献者的架构文档。Runtime 整体架构、模块设计、核心组件解析 \| [架构指南](design/README.md) \|
21	-<thead>	21	+\| 📝 研发规范 \| 面向贡献者的规范指南，包括设计文档模板、编码规范、测试规范、代码检视规则 \| [研发规范与贡献指南](guidelines/README.md) \|
22	-<tr>
23	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">文档</th>
24	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">定位与内容</th>
25	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">入口</th>
26	-</tr>
27	-</thead>
28	-<tbody>
29	-<tr>
30	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>⚡ 快速入门</b></td>
31	-<td style="word-wrap: break-word; overflow-wrap: break-word;">零基础起步，第一次接触 Runtime 的开发者。Runtime 简介 + 编程模型讲解，建议顺序阅读</td>
32	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="01_quick_start/Runtime简介.md">Runtime 简介</a> · <a href="01_quick_start/Runtime编程模型.md">编程模型</a></td>
33	-</tr>
34	-<tr>
35	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>📖 编程指南</b></td>
36	-<td style="word-wrap: break-word; overflow-wrap: break-word;">深度开发手册，已跑通入门阶段 Hello CANN 后需要深入理解原理的开发者。原理解析 + 示例代码</td>
37	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="02_dev_guide/00_dev_guide.md">编程指南</a></td>
38	-</tr>
39	-<tr>
40	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>📋 API 参考</b></td>
41	-<td style="word-wrap: break-word; overflow-wrap: break-word;">接口字典，开发过程中随时查阅函数定义。函数签名 + 参数说明 + 返回值</td>
42	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="03_api_ref/01_概述.md">头文件说明</a> · <a href="03_api_ref/api_ref.md">API 参考</a></td>
43	-</tr>
44	-<tr>
45	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>🏗️ 架构指南</b></td>
46	-<td style="word-wrap: break-word; overflow-wrap: break-word;">面向贡献者的架构文档。Runtime 整体架构、模块设计、核心组件解析</td>
47	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="design/README.md">架构指南</a></td>
48	-</tr>
49	-<tr>
50	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>📝 研发规范</b></td>
51	-<td style="word-wrap: break-word; overflow-wrap: break-word;">面向贡献者的规范指南，包括设计文档模板、编码规范、测试规范、代码检视规则</td>
52	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="guidelines/README.md">研发规范与贡献指南</a></td>
53	-</tr>
54	-</tbody>
55	-</table>
56		22
57	---	23	---
58		24
	@@ -63,156 +29,44 @@
63		29
64	### 🌱 入门阶段	30	### 🌱 入门阶段
65		31
66	-<table width="100%" style="table-layout: fixed; width: 100%; border-collapse: collapse;">	32	+\| 学习步骤 \| 内容说明 \| 入口 \|
67	-<colgroup>	33	+\|---------\|---------\|------\|
68	-<col width="15%">	34	+\| 环境准备 \| CANN 一键安装 \| [CANN 一键安装](https://www.hiascend.com/cann/download) \|
69	-<col width="55%">	35	+\| 概念原理 \| Runtime 核心概念与编程模型：Host-Device 架构、Context、Stream、同步异步、典型执行流程 \| [Runtime 简介](01_quick_start/Runtime简介.md) · [编程模型](01_quick_start/Runtime编程模型.md) \|
70	-<col width="30%">	36	+\| Hello CANN \| 第一个可运行的 Runtime 程序，完成最小计算闭环 \| [Hello CANN](../example/0_quickstart/0_hello_cann/README.md) \|
71	-</colgroup>
72	-<thead>
73	-<tr>
74	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">学习步骤</th>
75	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">内容说明</th>
76	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">入口</th>
77	-</tr>
78	-</thead>
79	-<tbody>
80	-<tr>
81	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>环境准备</b></td>
82	-<td style="word-wrap: break-word; overflow-wrap: break-word;">CANN 一键安装</td>
83	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="https://www.hiascend.com/cann/download">CANN 一键安装</a></td>
84	-</tr>
85	-<tr>
86	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>概念原理</b></td>
87	-<td style="word-wrap: break-word; overflow-wrap: break-word;">Runtime 核心概念与编程模型：Host-Device 架构、Context、Stream、同步异步、典型执行流程</td>
88	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="01_quick_start/Runtime简介.md">Runtime 简介</a> · <a href="01_quick_start/Runtime编程模型.md">编程模型</a></td>
89	-</tr>
90	-<tr>
91	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>Hello CANN</b></td>
92	-<td style="word-wrap: break-word; overflow-wrap: break-word;">第一个可运行的 Runtime 程序，完成最小计算闭环</td>
93	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/0_quickstart/0_hello_cann/README.md">Hello CANN</a></td>
94	-</tr>
95	-</tbody>
96	-</table>
97		37
98	---	38	---
99		39
100	### 🚀 进阶阶段	40	### 🚀 进阶阶段
101		41
102	-<table width="100%" style="table-layout: fixed; width: 100%; border-collapse: collapse;">	42	+\| 主题 \| 核心内容 \| 对应样例 \|
103	-<colgroup>	43	+\|------\|---------\|---------\|
104	-<col width="15%">	44	+\| 初始化 \| 包括环境初始化、设备资源配置、日志管理等 \| [device_normal](../example/1_basic_features/device/0_device_normal/README.md) \|
105	-<col width="55%">	45	+\| 内存管理 \| 包括 Host 内存管理、Device 内存管理、多流同步内存、内存拷贝（同步异步）、物理内存共享 (pid) 等 \| [h2d_sync_memory_copy](../example/1_basic_features/memory/1_h2d_sync_memory_copy/README.md) \|
106	-<col width="30%">	46	+\| 异步任务 \| 包括 Stream 管理、Event 管理、Kernel 加载与执行、内存语义同步等 \| [simple_stream](../example/1_basic_features/stream/0_simple_stream/README.md) \|
107	-</colgroup>
108	-<thead>
109	-<tr>
110	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">主题</th>
111	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">核心内容</th>
112	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">对应样例</th>
113	-</tr>
114	-</thead>
115	-<tbody>
116	-<tr>
117	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>初始化</b></td>
118	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括环境初始化、设备资源配置、日志管理等</td>
119	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/1_basic_features/device/0_device_normal/README.md">device_normal</a></td>
120	-</tr>
121	-<tr>
122	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>内存管理</b></td>
123	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括 Host 内存管理、Device 内存管理、多流同步内存、内存拷贝（同步/异步）、物理内存共享 (pid) 等</td>
124	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/1_basic_features/memory/1_h2d_sync_memory_copy/README.md">h2d_sync_memory_copy</a></td>
125	-</tr>
126	-<tr>
127	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>异步任务</b></td>
128	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括 Stream 管理、Event 管理、Kernel 加载与执行、内存语义同步等</td>
129	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/1_basic_features/stream/0_simple_stream/README.md">simple_stream</a></td>
130	-</tr>
131	-</tbody>
132	-</table>
133		47
134	---	48	---
135		49
136	### 🔥 高级阶段	50	### 🔥 高级阶段
137		51
138	-<table width="100%" style="table-layout: fixed; width: 100%; border-collapse: collapse;">	52	+\| 主题 \| 核心内容 \| 对应样例 \|
139	-<colgroup>	53	+\|------\|---------\|---------\|
140	-<col width="15%">	54	+\| ACL Graph \| 包括单流捕获、跨流捕获、任务更新等 \| [model_update](../example/2_advanced_features/model_ri/1_model_update/README.md) \|
141	-<col width="55%">	55	+\| 多设备编程 \| 包括跨 Device 数据交互、P2P 内存访问、多卡并行调度等 \| [device_P2P](../example/1_basic_features/device/2_device_P2P/README.md) \|
142	-<col width="30%">	56	+\| 进程间通信 \| 包括 IPC Event 同步、IPC 内存共享（指定 PID/不指定 PID）等 \| [ipc_event](../example/2_advanced_features/ipcevent/0_ipcevent/README.md) · [ipc_memory](../example/1_basic_features/memory/11_ipc_memory_withoutpid/README.md) \|
143	-</colgroup>	57	+\| 性能调优 \| 包括 Profiling 采集并落盘、获取网络模型中算子的性能数据、可视化展示原始性能数据解析结果等 \| [create_config](../example/5_performance/profiling/0_create_config/README.md) \|
144	-<thead>
145	-<tr>
146	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">主题</th>
147	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">核心内容</th>
148	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">对应样例</th>
149	-</tr>
150	-</thead>
151	-<tbody>
152	-<tr>
153	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>ACL Graph</b></td>
154	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括单流捕获、跨流捕获、任务更新等</td>
155	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/2_advanced_features/model_ri/1_model_update/README.md">model_update</a></td>
156	-</tr>
157	-<tr>
158	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>多设备编程</b></td>
159	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括跨 Device 数据交互、P2P 内存访问、多卡并行调度等</td>
160	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/1_basic_features/device/2_device_P2P/README.md">device_P2P</a></td>
161	-</tr>
162	-<tr>
163	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>进程间通信</b></td>
164	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括 IPC Event 同步、IPC 内存共享（指定 PID/不指定 PID）等</td>
165	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/2_advanced_features/ipcevent/0_ipcevent/README.md">ipc_event</a> · <a href="../example/1_basic_features/memory/11_ipc_memory_withoutpid/README.md">ipc_memory</a></td>
166	-</tr>
167	-<tr>
168	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>性能调优</b></td>
169	-<td style="word-wrap: break-word; overflow-wrap: break-word;">包括 Profiling 采集并落盘、获取网络模型中算子的性能数据、可视化展示原始性能数据解析结果等</td>
170	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="../example/5_performance/profiling/0_create_config/README.md">create_config</a></td>
171	-</tr>
172	-</tbody>
173	-</table>
174		58
175	---	59	---
176		60
177	<!-- ==================== 常见问题 ==================== -->	61	<!-- ==================== 常见问题 ==================== -->
178	## ❓ 常见问题	62	## ❓ 常见问题
179		63
180	-<table width="100%" style="table-layout: fixed; width: 100%; border-collapse: collapse;">	64	+\| 问题类型 \| 典型场景 \| 入口 \|
181	-<colgroup>	65	+\|---------\|---------\|------\|
182	-<col width="15%">	66	+\| 入门阶段 \| 初始化失败、Device 配置、版本兼容等入门常见问题 \| [版本不匹配](04_FAQ/Runtime版本与CANN版本不匹配导致的问题.md) · [aclInit 失败](04_FAQ/aclInit初始化失败常见原因排查.md) · [aclrtSetDevice 失败](04_FAQ/aclrtSetDevice调用失败.md) · [默认机制](04_FAQ/如何理解默认Device和默认Stream机制.md) \|
183	-<col width="55%">	67	+\| 基础开发 \| 内存管理、Stream 同步、数据复制等基础 API 使用问题 \| [内存申请失败](04_FAQ/aclrtMalloc内存申请失败常见原因.md) · [Stream 下发失败](04_FAQ/aclrtMemcpyAsync在错误的Stream上下发失败.md) · [内存策略](04_FAQ/如何选择合适的内存分配策略.md) · [同步机制](04_FAQ/Stream同步与Event同步的区别与选择.md) \|
184	-<col width="30%">	68	+\| 进阶场景 \| 多设备编程、ACL Graph、进程通信等复杂场景问题 \| [多Device Stream](04_FAQ/多Device场景下Stream跨Device下发失败.md) · [ACL Graph 任务提交](04_FAQ/ACLGraph捕获过程中任务提交限制.md) · [IPC 页表对齐](04_FAQ/进程间IPC内存共享的页表对齐要求.md) · [P2P 配置失败](04_FAQ/跨DeviceP2P数据交互配置失败.md) \|
185	-</colgroup>	69	+\| 错误排查 \| 错误码解读、算子异常、日志定位等诊断方法 \| [异步错误码](04_FAQ/如何获取和解读Runtime异步错误码.md) · [算子输出异常](04_FAQ/算子执行输出全0的常见原因排查.md) · [遇错即停定位](04_FAQ/遇错即停模式下错误定位方法.md) · [plog 日志定位](04_FAQ/如何通过plog日志定位Device侧异常.md) \|
186	-<thead>
187	-<tr>
188	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">问题类型</th>
189	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">典型场景</th>
190	-<th align="left" style="word-wrap: break-word; overflow-wrap: break-word;">入口</th>
191	-</tr>
192	-</thead>
193	-<tbody>
194	-<tr>
195	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>进程重启失败</b></td>
196	-<td style="word-wrap: break-word; overflow-wrap: break-word;">进程重启失败、资源申请不到</td>
197	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="04_FAQ/用户进程异常退出后重启进程失败.md">用户进程异常退出后重启进程失败</a></td>
198	-</tr>
199	-<tr>
200	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>内核版本不兼容</b></td>
201	-<td style="word-wrap: break-word; overflow-wrap: break-word;">算子输出全为 0、内核版本不兼容</td>
202	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="04_FAQ/低版本内核使用asan导致算子执行失败.md">低版本内核使用 asan 导致算子执行失败</a></td>
203	-</tr>
204	-<tr>
205	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>休眠唤醒失败</b></td>
206	-<td style="word-wrap: break-word; overflow-wrap: break-word;">休眠失败、hwts busy</td>
207	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="04_FAQ/AI应用进程未退出导致休眠唤醒失败.md">AI 应用进程未退出导致休眠唤醒失败</a></td>
208	-</tr>
209	-<tr>
210	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><b>析构崩溃</b></td>
211	-<td style="word-wrap: break-word; overflow-wrap: break-word;">应用程序运行过程中出现 coredump 导致了崩溃，应用程序异常终止</td>
212	-<td style="word-wrap: break-word; overflow-wrap: break-word;"><a href="04_FAQ/析构函数中调用aclFinalize导致应用进程coredump.md">析构函数中调用 aclFinalize 导致应用进程 coredump</a></td>
213	-</tr>
214	-</tbody>
215	-</table>
216		70
217	---	71	---
218		72