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

[tde_m]

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

[HRZ]

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]

I see,

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

Thanks

[Hiroto_K_Intel]

Besically, I agree with HRZ's comment and I'd like to follow up.

According to below latest document(Including Structure Data Types as Arguments in OpenCL Kernels), you can see it.

https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807965224.html#mwh1391807954819

If you have any more error, please let me know.

[tde_m]

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]

Hi, tde_m,

Yes, you are right.

BTW, could you like to test with -fast-emulator flag ?

Intel recommends when emulating kernel code, you add -fast-emulator for 18.1 and later.

https://www.intel.com/content/www/us/en/programmable/documentation/mwh1391807965224.html#hzx1553406416602

On my side, the code you posted firstly completed compiling successfully.

Only -march=emulator mode is legacy now and it is possibility that the issue was reported and has been resolved when migrating from lagacy emulator to fast emulator.

If so, I could not get the answer why, sorry.

[tde_m]

Hello,

I've tried but the aoc compiler returns with an error message:

Error: aoc: The Intel(R) Kernel Builder for OpenCL(TM) compiler (ioc64) can not be found

I'm using Version 18.1.1 Build 263 Pro Edition.

I've tried also on a different installation and version of the tool (19.1 Pro) and the error is always the same.

[Hiroto_K_Intel]

Hello,

ioc64 is original component in system studio.

Did you install intel_sdk_for_opencl_setup_<version>.run ?

It is included in AOCLPro tar package(when uncompressed, it is in components) and when installing by using setup.sh, this run script should be executed.

(If X11 windows disabled, the installation might fail)

If not installed, could you please install it and try again ?

If installed, please check the directory and you can see ioc64 binary.

the default path as non-root user is here.

/home/<user_name>/intel/system_studio_2019/opencl/SDK/bin

[Hiroto_K_Intel]

Hi, tde_m,

Following previous comment, did you resolve this issue ? If not, please let me know.

[tde_m]

Hello,

I've installed it using the provided .run file. However, since I've installed it on a server, I think that it has been disabled because X11 is not present.

Is there a way to install it without X-server?

Thanks

[Hiroto_K_Intel]

Hi, tde_m,

Without X11, you have to run some scripts with options manually.

./intel_sdk_for_opencl_setup_2019.3.run --nox11 --target ./<exacted_installer_dir> cd <exacted_installer_dir> #edit silent.cfg -> ACCEPT_EULA=accept #edit silent.cfg -> PSET_INSTALL_DIR=<package_install_target> ./install.sh -s silent.cfg

If you have questions about silent.cfg, please read README.txt in <exacted_installer_dir>

[Hiroto_K_Intel]

Hi, tde_m,

Did you resolve this issue ? or have any more question ?

[Hiroto_K_Intel]

Hi, tde_m,

Did you compile emulation successfully ? If not, could you like to explain error messages ?​

[tde_m]

Hello,

sorry for the delay.

I am working on version 19.1 of Quartus Pro. After extraction, I don't find the file that you are mentioning.

I tried to re-install in text-mode (setting the mode flag to "text"), but it still don't find the ioc64 binary.

Do I need to install the last version of Quartus?

Regards,

[Hiroto_K_Intel]

Hi, tde_m,

Did you compile your project successfully ?​ If you have any more issue, please let me know.

[Hiroto_K_Intel]

Hi, tde_m,

Do you have any roadblock about this issue ? If so, could you like to tell me it ?

[tde_m]

Hi Hiroto,

as mentioned (see my reply three messages above), I am not able to find the installation file that you were mentioning.

I just gave up on this, avoiding passing the user defined structure to the kernel.

[Hiroto_K_Intel]

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