当前位置：首页 > news >正文

CANN/runtime遇错即停错误定位

news 2026/7/11 11:33:35

遇错即停模式下错误定位方法

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime

问题现象描述

现象1：遇错即停模式下无法准确定位失败任务

在Stream上配置了遇错即停模式（ACL_STOP_ON_FAILURE），任务执行失败时整个Context下的任务被停止，但无法准确定位是哪个任务失败，失败原因是什么。

典型场景：

aclrtStream stream; aclrtCreateStream(&stream); aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE); // 下发多个任务 aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream); kernel1<<<8, nullptr, stream>>>(args1); kernel2<<<8, nullptr, stream>>>(args2); aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream); // 同步时发现错误，但不知道是哪个任务失败 aclError error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { // 如何定位是kernel1还是kernel2失败？ // 失败原因是什么？ }

可能原因

任务序列较长：Stream上下发了多个异步任务，错误发生后无法直观判断是哪个任务失败。
错误信息不够详细：仅通过返回码无法获知具体失败原因，需要结合plog日志和错误查询接口。
遇错即停特性限制：遇错即停模式会在首个错误发生时停止Context中所有Stream的任务执行，可能影响错误现场的保护。

处理步骤

方法1：分段同步定位错误任务

将长任务序列分解为多个小段，每段后进行同步检查：

aclrtStream stream; aclrtCreateStream(&stream); aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE); // 任务1：Host -> Device拷贝 aclError error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream); if (error != ACL_RT_SUCCESS) { printf("Task 1 launch failed: %d\n", error); return error; } error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { printf("Task 1 execution failed: %d\n", error); char *errMsg = aclGetRecentErrMsg(); printf("Error: %s\n", errMsg); return error; } // 任务2：kernel1执行 kernel1<<<8, nullptr, stream>>>(args1); error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { printf("Task 2 (kernel1) execution failed: %d\n", error); char *errMsg = aclGetRecentErrMsg(); printf("Error: %s\n", errMsg); return error; } // 任务3：kernel2执行 kernel2<<<8, nullptr, stream>>>(args2); error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { printf("Task 3 (kernel2) execution failed: %d\n", error); char *errMsg = aclGetRecentErrMsg(); printf("Error: %s\n", errMsg); return error; } // 任务4：Device -> Host拷贝 aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream); error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { printf("Task 4 execution failed: %d\n", error); return error; } aclrtDestroyStream(stream);

方法2：使用Event标记任务边界

在关键任务之间插入Event，通过Event状态判断任务执行进度：

aclrtStream stream; aclrtCreateStream(&stream); aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE); // 创建多个Event用于标记任务边界 aclrtEvent event1, event2, event3; aclrtCreateEvent(&event1); aclrtCreateEvent(&event2); aclrtCreateEvent(&event3); // 下发任务并插入Event标记 aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream); aclrtRecordEvent(event1, stream); // 标记拷贝任务完成 kernel1<<<8, nullptr, stream>>>(args1); aclrtRecordEvent(event2, stream); // 标记kernel1完成 kernel2<<<8, nullptr, stream>>>(args2); aclrtRecordEvent(event3, stream); // 标记kernel2完成 aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream); // 同步并检查错误 aclError error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { // 查询Event状态，定位失败任务 aclrtEventStatus status1, status2, status3; aclrtEventQuery(event1, &status1); aclrtEventQuery(event2, &status2); aclrtEventQuery(event3, &status3); printf("Event status: event1=%d, event2=%d, event3=%d\n", status1, status2, status3); // ACL_RT_EVENT_COMPLETE = 1, ACL_RT_EVENT_NOT_READY = 0 // 最后一个完成的Event之后的任务为失败任务 if (status1 == ACL_RT_EVENT_COMPLETE && status2 == ACL_RT_EVENT_NOT_READY) { printf("Task failed between event1 and event2 (kernel1 failed)\n"); } else if (status2 == ACL_RT_EVENT_COMPLETE && status3 == ACL_RT_EVENT_NOT_READY) { printf("Task failed between event2 and event3 (kernel2 failed)\n"); } char *errMsg = aclGetRecentErrMsg(); printf("Error detail: %s\n", errMsg); } // 清理资源 aclrtDestroyEvent(event1); aclrtDestroyEvent(event2); aclrtDestroyEvent(event3); aclrtDestroyStream(stream);

方法3：结合错误查询接口和plog日志

通过错误查询接口获取错误码，结合plog日志定位具体失败原因：

aclrtStream stream; aclrtCreateStream(&stream); aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE); // 下发任务 // ...（省略任务下发代码） // 同步并检查错误 aclError error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { // 1. 获取错误码 printf("Stream synchronize failed with error code: %d (0x%x)\n", error, error); // 2. 查看最近错误 aclError lastError = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL); printf("Last error: %d (0x%x)\n", lastError, lastError); // 3. 获取详细错误信息 char *errMsg = aclGetRecentErrMsg(); if (errMsg != nullptr) { printf("Error message: %s\n", errMsg); } // 4. 提示查看plog日志 printf("Please check plog for detailed error information:\n"); printf(" - Look for 'fault kernel_name' to identify failed kernel\n"); printf(" - Look for 'retCode' to get error code\n"); printf(" - Look for 'error module_type' and 'module_name' for error source\n"); } aclrtDestroyStream(stream);

plog日志关键信息提取：

# 查找失败的任务 grep "Task run failed" plog.log # 查找失败的内核名称 grep "fault kernel_name" plog.log # 查找错误码和类型 grep "retCode\|errType" plog.log # 查找错误模块 grep "module_type\|module_name" plog.log # 示例输出解析： # [ERROR] Task run failed, device_id=0, stream_id=2, task_id=1 # -> 定位到具体的设备、流、任务ID # fault kernel_name=Add_ee98c6628030785f610b924ab1557b31 # -> 定位到具体失败的内核 # retCode=0x31, [vector core exception] # -> 错误原因：向量核异常 # module_type=5, module_name=EZ9999 # -> 错误来源模块

方法4：对比遇错继续模式

如果遇错即停模式难以定位，可临时切换到遇错继续模式，让后续任务继续执行以收集更多信息：

aclrtStream stream; aclrtCreateStream(&stream); // 测试时使用遇错继续模式 aclrtSetStreamFailureMode(stream, ACL_CONTINUE_ON_FAILURE); // 下发多个任务 // ... // 同步并检查错误 aclError error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { printf("Error occurred: %d\n", error); // 遇错继续模式下，可以获取更多信息 // 但注意：aclrtPeekAtLastError可能返回的是最后发生的错误，而非首次错误 aclError lastError = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL); printf("Last error: %d\n", lastError); char *errMsg = aclGetRecentErrMsg(); printf("Error message: %s\n", errMsg); // 查看哪些任务执行成功，哪些失败 // （需要通过中间结果验证） } // 问题定位后，改回遇错即停模式用于生产环境 aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE); aclrtDestroyStream(stream);

方法5：使用Host回调任务进行调试

在关键任务之间插入Host回调函数，记录任务执行进度：

void taskCallback(void *userData) { int *taskId = (int*)userData; printf("Task %d completed successfully\n", *taskId); } // ... aclrtStream stream; aclrtCreateStream(&stream); aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE); int taskId1 = 1, taskId2 = 2, taskId3 = 3, taskId4 = 4; // 下发任务并插入回调 aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream); aclrtLaunchHostFunc(stream, taskCallback, &taskId1); kernel1<<<8, nullptr, stream>>>(args1); aclrtLaunchHostFunc(stream, taskCallback, &taskId2); kernel2<<<8, nullptr, stream>>>(args2); aclrtLaunchHostFunc(stream, taskCallback, &taskId3); aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream); aclrtLaunchHostFunc(stream, taskCallback, &taskId4); // 同步 aclError error = aclrtSynchronizeStream(stream); if (error != ACL_RT_SUCCESS) { // 通过回调输出可以知道最后一个成功完成的任务 // 失败的任务就是最后一个成功任务之后的任务 printf("Error occurred after last successful task\n"); } aclrtDestroyStream(stream);