Forum Discussion

tde_m's avatar
tde_m
Icon for Occasional Contributor rankOccasional Contributor
6 years ago

Emulation fails with user-defined data type for a kernel argument

Hello,

as far as I understand, it is plausible to use a user-defined data type for a kernel argument.

Example:

typedef struct{
    char my_id;
    char num_ids;
}my_type_t;
 
__kernel void app(const int N, my_type_t id, __global int * mem)
{
    for(int i=0;i<N;i++)
        mem[i]=id.my_id;
}

However, if I try to compile this for emulation I obtain an internal compiler error:

aoc -march=emulator  -board=p520_max_sg280l  bugs/kernel_arguments.cl -o bin/scaling_rank0.aocx   -fp-relaxed -fpc
aoc: Running OpenCL parser....
aoc: OpenCL parser completed successfully.
aoc: Linking Object files....
aoc: Compiling for Emulation ....
sh: line 1: 317621 Aborted                 /users/tdematt/intelFPGA_pro/18.1/hld/linux64/bin/../../llvm/bin/aocl-opt -translate-library-calls -reverse-library-translation -insert-ip-library-calls -create-emulator-wrapper -generate-emulator-sys-desc -emulDirCleanup -fp-relaxed=true -fpc=true -dbg-info-enabled "kernel_arguments.1.bc" -o "kernel_arguments.bc" >> kernel_arguments.log 2> opt.err
 
******* Error: Assert failure at /build/swbuild/SJ/nightly/18.1.1/263/l64/p4/acl/llvm/lib/Analysis/FPGAAnalysis/GenerateEmulatorSysDesc.cpp(328) *******
!"Unknown arg memory space" FAILED
/users/tdematt/intelFPGA_pro/18.1/hld/llvm/bin/../lib/libLLVM-6.0.so(_ZN4llvm3sys15PrintStackTraceERNS_11raw_ostreamE+0x2a)[0x7f79bdaa377a]
/users/tdematt/intelFPGA_pro/18.1/hld/llvm/bin/../lib/libLLVM-6.0.so(_ZN4llvm3sys17RunSignalHandlersEv+0x3e)[0x7f79bdaa14be]
/users/tdematt/intelFPGA_pro/18.1/hld/llvm/bin/../lib/libLLVM-6.0.so(+0xb0d864)[0x7f79bdaa1864]
/lib64/libpthread.so.0(+0xf5d0)[0x7f79bcd875d0]
/lib64/libc.so.6(gsignal+0x37)[0x7f79bc146207]
/lib64/libc.so.6(abort+0x148)[0x7f79bc1478f8]
/users/tdematt/intelFPGA_pro/18.1/hld/llvm/bin/../lib/libLLVM-6.0.so(_ZN3acl23GenerateEmulatorSysDesc11runOnModuleERN4llvm6ModuleE+0x15cc)[0x7f79bf183aec]
/users/tdematt/intelFPGA_pro/18.1/hld/llvm/bin/../lib/libLLVM-6.0.so(_ZN4llvm6legacy15PassManagerImpl3runERNS_6ModuleE+0x2d8)[0x7f79bdbb7718]
/users/tdematt/intelFPGA_pro/18.1/hld/linux64/bin/../../llvm/bin/aocl-opt(main+0x236b)[0x55bf2d78f21b]
/lib64/libc.so.6(__libc_start_main+0xf5)[0x7f79bc1323d5]
/users/tdematt/intelFPGA_pro/18.1/hld/linux64/bin/../../llvm/bin/aocl-opt(+0x27fad)[0x55bf2d78ffad]
Stack dump:
0.      Program arguments: /users/tdematt/intelFPGA_pro/18.1/hld/linux64/bin/../../llvm/bin/aocl-opt -translate-library-calls -reverse-library-translation -insert-ip-library-calls -create-emulator-wrapper -generate-emulator-sys-desc -emulDirCleanup -fp-relaxed=true -fpc=true -dbg-info-enabled kernel_arguments.1.bc -o kernel_arguments.bc
1.      Running pass 'GenerateEmulatorSysDesc: Creates kernel wrapper functions and lookup table that can be called from the emulator run time' on module 'kernel_arguments.1.bc'.
Error: Optimizer FAILED.
Refer to scaling_rank0/kernel_arguments.log for details.

I was wondering if this is still a legit piece of code also for Intel OpenCL.

The code was compiled against a stratix 10, using bot v18.1 and 19.1 versions of the compiler.

Thanks

17 Replies

  • HRZ's avatar
    HRZ
    Icon for Frequent Contributor rankFrequent Contributor
    6 years ago

    I don't think you can use anything other than OpenCL's datatypes as kernel argument, unless the argument is a pointer to global memory. Even though your code compiles just fine with recent versions of the compiler and just crashes during emulation, with older versions of the compiler (below v18.1) a clear error is printed as follows:

    error: unsupported kernel argument type
__kernel void app(const int N, my_type_t id, __global int * mem)
                                         ^

    The following code works fine with every version of the compiler and in both normal compilation and emulation:

    typedef struct{
   char my_id;
   char num_ids;
}my_type_t;
 
__kernel void app(const int N, __global my_type_t* id, __global int * mem)
{
   for(int i=0;i<N;i++)
	  mem[i]=id->my_id;
}

    I think the previous behavior of the compiler (below v18.1) was correct and the fact that your original code compiles fine starting from v18.1 (but crashes during emulation) is the incorrect behavior.

    P.S. You should probably consider using the __attribute__ ((packed)) to avoid padding differences between host and kernel for the struct.

    • tde_m's avatar
      tde_m
      Icon for Occasional Contributor rankOccasional Contributor
      6 years ago

      I see,

      thanks I will pass through device memory (or use standard opencl types).

      Thanks

    • tde_m's avatar
      tde_m
      Icon for Occasional Contributor rankOccasional Contributor
      6 years ago

      Hi Hiroto,

      so according to the documentation I can pass structure parameters in OpenCL kernels by value.

      But the emulator fails in compiling it (see my first post). Why is this happening?

      This occur even if I disable the automatic alignment by using the packed attribute (in any case, the alignment is 1 byte being the struct composed by 2 chars).

      Thanks

  • Hiroto_K_Intel's avatar
    Hiroto_K_Intel
    Icon for New Contributor rankNew Contributor
    6 years ago

    Hi, tde_m,

    I missed your messages because of the system bugs.

    I'm really sorry for your inconvenience.

    Then, I can read your message now and I understood your situation.

    ​

    Regarding installations, did you download Intel FPGA SDK for OpenCL packages ? If so, you could get AOCL-pro-<version>-linux.tar. This link is the detailed answer and I'd like to read it.

    https://www.intel.com/content/altera-www/global/en_us/index/support/support-resources/knowledge-base/embedded/2019/error--aoc--the-intel-r--kernel-builder-for-opencl-tm--compiler-0.html

    ​