Forum Discussion

Altera_Forum's avatar
Altera_Forum
Icon for Honored Contributor rankHonored Contributor
7 years ago

Error: Assert failure at ...MemBankPortAssign.cpp(191)

Hi,

When I tried compiler command aoc -c ... --report for my kernel, it stuck after showing the warnings... I did not know why so I then tried to comment parts of my kernel out to isolate the lines of code that cause the error.

However, when my kernel only has these lines： 

# pragma OPENCL EXTENSION cl_khr_fp64 : enable# pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable# pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
__kernel void test(int A, int B, int b, __local int *gid_, __global double *input,
    __global int *finished, __global int *head) {
    const int tid = get_local_id(0);
    int       m   = A * B - 1;
    if(tid == 0) // Dynamic fetch
        gid_         = atomic_add(&head, 1);
    barrier(CLK_LOCAL_MEM_FENCE);
    while(gid_ < m) {
        if(tid == 0) // Dynamic fetch
            gid_     = atomic_add(&head, 1);
        barrier(CLK_LOCAL_MEM_FENCE);
    }
}

I cannot generate the report as I get the "Error: Assert failure at ...MemBankPortAssign.cpp(191)". 


******* Error: Assert failure at d:/SJ/nightly/16.1/196/w64/p4/acl/llvm/lib/Transforms/FPGATransforms/MemBankPortAssign.cpp(191) *******
m_num_ports_per_bank > 0 FAILED
Stack dump:
0.	Program arguments: C:/intelFPGA_pro/16.1/hld/windows64/bin/aocl-opt --acle ljg7wk8o12ectgfpthjmnj8xmgf1qb17frkzwewi22etqs0o0cvorlvczrk7mipp8xd3egwiyx713svzw3kmlt8clxdbqoypaxbbyw0oygu1nsyzek
......
wcmgd33czpygfb0gwieguzqgy7atj3qe8clgssmxw7ljg70rji72eqqr0oekh3njjx2rjbtijzgxf38uui3xyhmujp7tkelkybygpsmvwo12hswkvm7jlzmtyz23gbnf0318a7qxwow2ko7u8zbgwknjf0a -board c:/intelFPGA_pro/16.1/hld/board/de5a_net_e1/hardware/de5a_net_e1/board_spec.xml -dbg-info-enabled test.1.bc -o test.kwgid.bc 
1.	Running pass 'Pick memory configuration' on module 'test.1.bc'.
2.	Running pass 'Memory Bank/Port Assignments.' on function '@PTTWAC_soa_asta'
0x000000013FB6CE37 (0x0000000000B3EF30 0x00000000032FBC80 0x0000000003DE4DF0 0x0000000000000001)
0x00000001402BA949 (0x00000000033F4DC0 0x00000000033F4DC0 0x00000000000000E0 0x0000000000000000)
0x00000001402BB4BD (0x00000000032FBC80 0x0000000000B3F2F0 0x0000000000000004 0x0000000003C81D50)
0x00000001402BB87C (0x0000000003382880 0x0000000000000000 0x0000000000000000 0x00000001409CF048)
0x000000013FB096C3 (0x0000000000000000 0x00000000013BA460 0x000000000000000C 0x000000000336E6D0)
0x000000013FB158DC (0x0000000000000000 0x0000000000000000 0x00000000032FB050 0x0000000003CF58C0)
0x000000013FB1730D (0x0000000000000000 0x0000000003CF58C0 0x0000000003CF58C8 0x00000000032FB050)
0x000000013FB17B3B (0x0000000001321B40 0x00000000032FAEB0 0x0000000000000007 0x0000000000000000)
0x00000001402BAD6D (0x0000000000000035 0x000000000130B930 0x00000000000001A8 0x0000000000000000)
0x00000001402BB070 (0x0000000003C03F30 0x00000000000008F0 0x000000000000011E 0x0000000000B3F9C0)
0x000000013F7FED58 (0x0000000000000001 0x00000001A7FE8ADE 0x0000000000000001 0x0000000000000000)
0x00000001404407D6 (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), ??4_Init_locks@std@@QEAAAEAV01@AEBV01@@Z() + 0x7F6 bytes(s)
0x0000000076E759CD (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), BaseThreadInitThunk() + 0xD bytes(s)
0x0000000076FAA561 (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), RtlUserThreadStart() + 0x21 bytes(s)

Who can help me? Thanks in advance.

2 Replies

  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    7 years ago

    What do you mean by getting "stuck"? The OpenCL to LLVM conversion, after which the area estimation report is generated, could take up to an hour depending on the complexity of the kernel.

    The code snippet you have posted compiles fine on my machine using AOC 16.1.2. Note that the barriers in your code snippet are unnecessary since you do not have any local memory constructs in your kernel. Furthermore, atomic global memory operations will result in extremely poor performance.
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    7 years ago

    Thank you very much.

    You said you compiled fine, then I found the code I posted just removed one unused parameter (__local int *done). It was just this parameter caused this problem, although I did not know why. I have removed it from parameter list by just using it inside the kernel...

    Then with your help, I successfully commented parts of my kernel out and found the lines of code that cause the "stuck". I am now trying to solve the "stuck".