Forum Discussion

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

Error: Assert failure at ... acl::ACLMemorySpaces::is_local_addrspace((cast<A

I touch attribute "autorun" in HelloWorld project, write OpenCL program: 


global int i;
global int j;
__attribute__((max_global_work_dim(0)))
__attribute__((autorun))
__kernel void hello_world()
{
  i = 0;
  j = 1;
  while (1) ;
}

and fast see big ASSERT output in last Quartus version of aoc (16.1, 17.0): 


Compiler Command: aoc -v device/hello_world.cl -o bin/hello_world.aocx
******* Error: Assert failure at d:/SJ/nightly/17.0/290/w64/p4/acl/llvm/lib/FPGAAnalysis/ACLMemoryDependenceAnalysis2.cpp(3543) *******
acl::ACLMemorySpaces::is_local_addrspace((cast<ACLMemInst>(I))->getPointerAddressSpace()) FAILED
Stack dump:
0.      Program arguments: D:/intelFPGA_pro/17.0/hld/windows64/bin/aocl-opt --acle ljg7wk8o12ectgfpthjmnj8xmgf1qb17frkzwewi22etqs0o0cvorlvc
zrk7mipp8xd3egwiyx713svzw3kmlt8clxdbqoypaxbbyw0oygu1nsyzekh3nt0x0jpsmvypfxguwwdo880qqk8pachqllyc18a7q3wp12j7eqwipxw13swz1bp7tk71wyb3rb17frk
...
3ndpos3fcle0burjbtijz3gfuwwjz880qqkdoehdqlr8v0jpsmvpz3rj1wjdor2qmqw07akc -board c:/intelFPGA_pro/16.1/hld/board/alaric_hpc_16.1.2_b203_2ban
ks_256_stp/hardware/alaric_v3_prod_2banks_256/board_spec.xml -dbg-info-enabled --grif --soft-elementary-math=false --fas=false --wiicm-disa
ble=true hello_world.1.bc -o hello_world.kwgid.bc
1.      Running pass 'Function Pass Manager' on module 'hello_world.1.bc'.
2.      Running pass 'Pipeline Memory Dependence Analysis' on function '@hello_world'
0x000000013FEA8740 (0x00000000035F9870 0x0000000002BC9EE0 0x0000000002BDBF30 0x00000000035F9800)
0x0000000140364AB9 (0x0000000002BC9EE0 0x0000000002B6F750 0x0000000000000210 0x0000000000000000)
0x0000000140364C69 (0x0000000000000000 0x0000000000A6F901 0x0000000002B6F750 0x00000000034680C0)
0x0000000140364EDD (0x000000000000003A 0x0000000002F9E8E0 0x00000000000001D0 0x0000000000000000)
0x00000001403651E0 (0x000000000357F7F0 0x0000000000000998 0x0000000000000133 0x0000000000A6F940)
0x000000013F7FED58 (0x0000000000000001 0x0000001B3ED1F579 0x0000000000000001 0x0000000000000000)
0x00000001404EF0A6 (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), ??4_Init_locks@std@@QEAAAEAV01@AEBV01@@Z(
) + 0x80A bytes(s)
0x00000000776959CD (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), BaseThreadInitThunk() + 0xD bytes(s)
0x00000000778CA561 (0x0000000000000000 0x0000000000000000 0x0000000000000000 0x0000000000000000), RtlUserThreadStart() + 0x21 bytes(s)

If I declare and set to 0 only 1 of global variables "i" or "j", then .cl is compiled successfull, normal programmed to flash and runned on start of the my PC.

Where are problem ? In Altera recomendations aocl_programming_guide.pdf I can`t find a restrictions for amount of global variables and setting to its on fly.

In the standard opencl-1.0.pdf in section "6.5 Address Space Qualifiers" is quoted an example of using attribute "__global" for global variables: 


__global int *p;

These declaration in upper code is caused next compilation error: 


... /device/hello_world.cl:21:15: error: global variables must have a constant or channel address space qualifier
__global int *p;
              ^
1 error generated.
Error: OpenCL parser FAILED.
Refer to hello_world/hello_world.log for details.

Altera (and Intel) isn`t love and not test global variables in its version of Khronos OpenCL ?

3 Replies

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

    The answer to your problem is right in the same section in the OpenCL specification:

    --- Quote Start ---

    All program scope variables must be declared in the __constant address space.

    --- Quote End ---

    The specification allows to you declare __global pointers inside of a kernel (which can point to an existing __global address space passed onto the kernel as a __global pointer). You cannot have program scope __global variables that are shared between multiple kernels, simply beucase you cannot allocate memory inside the kernel or share data between kernels in OpenCL (unless you use explicit point-to-point communication via channels/pipes). You can, however, have program scope __constant variables that are initialized.

    As to why it works when you only declare one of the variables: it is likely because the variable is optimized out of the code. I'm not sure why the same thing does not happen when two variables are declared, though.

    Note that autorun kernels cannot have an interface to host or memory; if you are trying to read data from global memory in an autorun kernel, it is not going to work. You have to read it in a non-autorun kernel and pass it to the autorun kernel using channels.
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    8 years ago

    --- Quote Start ---

    Note that autorun kernels cannot have an interface to host or memory; if you are trying to read data from global memory in an autorun kernel, it is not going to work. You have to read it in a non-autorun kernel and pass it to the autorun kernel using channels.

    --- Quote End ---

    This is bad news. I want use "autorun" kernel with global memory opened to host for omit Altera/Intel hidden/secret interface to data exchange and make my "signals" and "replies" back in memory.

    In aocl_programming_guide.pdf (2.3.1.2) is no words about "no interface to memory", only 0 with host.

    Or Altera interface from host with OpenCL kernels is open ? Where I may read about it ?

    We want cut a working OpenCL core in VHDL and insert it to big FPGA project with more other no-OpenCL blocks, it is possible ?
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    8 years ago

    --- Quote Start ---

    This is bad news. I want use "autorun" kernel with global memory opened to host for omit Altera/Intel hidden/secret interface to data exchange and make my "signals" and "replies" back in memory.

    In aocl_programming_guide.pdf (2.3.1.2) is no words about "no interface to memory", only 0 with host.

    Or Altera interface from host with OpenCL kernels is open ? Where I may read about it ?

    We want cut a working OpenCL core in VHDL and insert it to big FPGA project with more other no-OpenCL blocks, it is possible ?

    --- Quote End ---

    You cannot allocate global memory from kernel, it has to be allocated on the host and the pointer passed onto the kernel. Since autorun kernels cannot have an interface to the host, they cannot receive the pointer to global memory and hence, direct communication between autorun kernels and external memory is impossible. Basically, an autorun kernel cannot have any arguments and it is not at all visible from the host. The only means of communication with an autorun kernel is through channels.

    On the subject of connecting an OpenCL kernel to a custom memory interface, that is not going to work unless your interface is 100% compatible with Altera's interface. I am not really sure whether their interface is open or not, but my guess is that it is not. For your specific usage, I think the new Intel HLS tool would be a much better choice since it allows you to convert a C/C++ component to an HDL module with a standard Avalon interface (Memory-mapped, streaming, etc.) and can be connected to any other HDL module using Qsys. You also have the choice of integrating your existing HDL modules inside of an OpenCL project, but the modules would need to comply with Altera's interface and communication with DDR/PCI-E/Network will have to go through Altera's interface.