[Bug-Report|缺陷反馈]: AscendC::InitGlobalMemory在输入size较大时出现了超时aicore-asc-devkit-AtomGit

已关闭

[Bug-Report|缺陷反馈]: AscendC::InitGlobalMemory在输入size较大时出现了超时aicore #1434

zzzyh22创建于 8月5日关闭于 8月25日

zzzyh22

8月5日创建

Thanks for sending an issue! Please fill in the following template to help quickly solve your problem.

Describe the current behavior / 问题描述 (Mandatory / 必填)

按照如下方式使用该API

在入参大小为980,591,780时出现aicore，注释该api后aicore消失

Environment / 环境信息 (Mandatory / 必填)

ASCEND950

Steps to reproduce the issue / 重现步骤 (Mandatory / 必填)

使用同样大小的size作为入参，调用该API

Describe the expected behavior / 预期结果 (Mandatory / 必填)

正常刷-INF

0x121801f08f51a,0x4240143018000000, mte error info: 0x43337a800020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 61, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x1218018c4f51a,0x4240143018000000, mte error info: 0x43337a800020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 62, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x121801baaf51a,0x4240143018000000, mte error info: 0x43337ac00020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 63, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x121800b86f51a,0x4240143018000000, mte error info: 0x43337ac00020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 64, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x1218020def51a,0x4240143018000000, mte error info: 0x43337b000020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 65, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x12180209cf51a,0x4240143018000000, mte error info: 0x43337b000020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 66, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x1218025fef51a,0x4240143018000000, mte error info: 0x43337b000020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E The error from device(chipId:0, dieId:0), serial number is 348, there is an aivec error exception, core id is 67, error code = 0, dump info: pc start: 0x12004102c000, current: 0x12004102de68, sc error info: 0xffffffffffff, su error info: 0x121801fc8f51a,0x4240143018000000, mte error info: 0x43337b400020799, vec error info: 0x410101d000393739, cube error info: 0, l1 error info: 0, aic error mask: 0x295856, para base: 0x12004c600000, mte error: 0, aic cond: 0.
E The extend info: errcode:(0) errorStr: timeout or trap error. subErrType: 0x4.
E For details, see the troubleshooting document on the Ascend official website. Search for the keyword "AI Core Error".[FUNC:ProcessDavidStarsCoreErrorInfo][FILE:device_error_proc_c.cc][LINE:749]
E An error occurred in the kernel task, retCode=0x25, [aicore timeout].[FUNC:PreCheckTaskErr][FILE:davinci_kernel_task.cc][LINE:1262]
E rtDeviceSynchronizeWithTimeout execution failed, reason=aicore timeout[FUNC:FuncErrorReason][FILE:error_message_manage.cc][LINE:69]
E wait for compute device to finish failed, runtime result = 507014.[FUNC:ReportCallError][FILE:log_inner.cpp][LINE:148]

Special notes for this issue/备注 (Optional / 选填)

lihuaichao成员

8月5日评论：

您好！

我们已收到此问题，将在确认后向您反馈~

ApeiriaNode_Booker成员

8月5日将 lihuaichao 设为负责人

lihuaichao成员

8月6日评论：

您好，感谢您的反馈。能否给出完整的代码或图示。

zzzyh22

8月6日评论：

您好，感谢您的反馈。能否给出完整的代码或图示。

@lihuaichao
使用代码如下，给出的大小为dealSize
if (constInfo_.returnValue) {
event_t eventIDMTE3ToV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_V));
SetFlagHardEvent::MTE3_V(eventIDMTE3ToV);
WaitFlagHardEvent::MTE3_V(eventIDMTE3ToV);

    GlobalTensor<uint16_t> valueOutGmTmp;
    valueOutGmTmp.SetGlobalBuffer((__gm__ uint16_t *)valueOutGm.GetPhyAddr());
    GlobalTensor<uint16_t> valueOut = valueOutGmTmp[paddingOffset];

    AscendC::InitGlobalMemory(valueOut, dealSize, constInfo_.NEG_INF_BFLOAT);
    SetFlag<HardEvent::MTE3_V>(eventIDMTE3ToV);
    WaitFlag<HardEvent::MTE3_V>(eventIDMTE3ToV);
}

lihuaichao成员

8月18日评论：

您好，InitGlobalMemory已废弃，对应新增接口为Fill，文档参考：https://gitcode.com/cann/asc-devkit/blob/master/docs/zh/api/SIMD-API/adv_api/tensor_transform/Fill.md

使用Fill接口，填充980591780个元素，未报错，功能正常。将如下代码保存至fill.asc后执行命令可复现。

bisheng fill.asc -o main --npu-arch=dav-3510 && ./main

#include <cstdint>
#include <cstdio>

#include "acl/acl.h"
#include "kernel_operator.h"

namespace MyCustomKernel {

constexpr uint64_t ELEMENT_COUNT = 980591780ULL;
constexpr float FILL_VALUE = 1.0F;
constexpr uint32_t NUM_BLOCKS = 1U;

class KernelFill {
public:
    __aicore__ inline void Init(__gm__ uint8_t* z)
    {
        zGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(z), ELEMENT_COUNT);
        AscendC::Fill(zGm, ELEMENT_COUNT, FILL_VALUE);
    }

private:
    AscendC::GlobalTensor<float> zGm;
};

} // namespace MyCustomKernel

extern "C" __global__ __vector__ void fill_custom(__gm__ uint8_t* z)
{
    AscendC::TPipe pipe;
    MyCustomKernel::KernelFill op;
    op.Init(z);
}

#define CHECK_ACL(expr)                                                                            \
    do {                                                                                           \
        const aclError ret = (expr);                                                              \
        if (ret != ACL_SUCCESS) {                                                                  \
            std::fprintf(stderr, "%s failed: %d\n", #expr, ret);                                \
            return 1;                                                                              \
        }                                                                                          \
    } while (0)

int main()
{
    constexpr size_t OUTPUT_BYTES = static_cast<size_t>(MyCustomKernel::ELEMENT_COUNT * sizeof(float));
    uint8_t* outputDevice = nullptr;
    float output = 0.0F;

    CHECK_ACL(aclInit(nullptr));
    CHECK_ACL(aclrtSetDevice(0));
    CHECK_ACL(aclrtMalloc(reinterpret_cast<void**>(&outputDevice), OUTPUT_BYTES, ACL_MEM_MALLOC_HUGE_FIRST));

    fill_custom<<<MyCustomKernel::NUM_BLOCKS, 0>>>(outputDevice);
    CHECK_ACL(aclrtSynchronizeDevice());
    CHECK_ACL(aclrtMemcpy(&output, sizeof(output), outputDevice, sizeof(output), ACL_MEMCPY_DEVICE_TO_HOST));

    CHECK_ACL(aclrtFree(outputDevice));
    CHECK_ACL(aclrtResetDevice(0));
    CHECK_ACL(aclFinalize());

    std::printf("first output element: %f\n", output);
    return output == MyCustomKernel::FILL_VALUE ? 0 : 1;
}

lihuaichao成员

8月25日评论：

您好，如果没有其他问题，我们先关闭此issue。后续如果有其他问题，欢迎再提issue交流~

L lihuaichao成员

8月25日 issue状态由待办的改变为已解决

L lihuaichao成员

8月25日关闭了 issue

L lihuaichao成员

21 天前添加了label：Accepted

Describe the current behavior / 问题描述 (Mandatory / 必填)

Environment / 环境信息 (Mandatory / 必填)

Steps to reproduce the issue / 重现步骤 (Mandatory / 必填)

Describe the expected behavior / 预期结果 (Mandatory / 必填)

Related log / screenshot / 日志 / 截图 (Mandatory / 必填)

Special notes for this issue/备注 (Optional / 选填)