"IOC Engine crashed" when building kernel

Chris_Juelg · ‎09-04-2013

typedef union _float3 {
    struct {float x, y, z;};
} _float3;

__kernel void myKernel(
        _float3 offset,
        __global float * output ) {

   output[0] += offset.z;
}

When using KernelBuilder64 to build this kernel, I receive "Error occured, IOC Engine crashed" in an error popup, no compiler output in the Console and the following on stdout:

Bitcast type must not be aggregate
%0 = bitcast i8* %pBuffer to %struct.anon
Instruction does not dominate all uses!
%0 = bitcast i8* %pBuffer to %struct.anon
%offset.coerce01.i = extractvalue %struct.anon %0, 0
Instruction does not dominate all uses!
%offset.coerce01.i = extractvalue %struct.anon %0, 0
%5 = fadd float %4, %offset.coerce01.i
Instruction does not dominate all uses!
%5 = fadd float %4, %offset.coerce01.i
store float %5, float addrspace(1)* %3, align 4
Broken module found, compilation aborted!

I use KernelBuilder64 3.0.0.1 from Intel® SDK for OpenCL* Applications XE 2013 on OpenSuse 12.2
The same kernel compiles without warning using the Nvidia OpenCL compiler. It also compiles as C99 with gcc when adapting the kernel function signature.

Just compiling this kernel in KernelBuilder works fine:

Setting target instruction set architecture to: Default
Intel OpenCL CPU device was found!
Device name: Intel(R) Core(TM) i7-3930K CPU @ 3.20GHz
Device version: OpenCL 1.2 (Build 67279)
Device vendor: Intel(R) Corporation
Device profile: FULL_PROFILE
Compilation started
Compilation done

Compilation succeeded!

Using a simple struct instead of the union of struct removes the error, but makes the rest of my program more complicated.

Chris_Juelg · ‎09-04-2013

i just noticed that the actual stdout message for the given kernel is:

Bitcast type must not be aggregate
%1 = bitcast i8* %0 to %union._float3
Instruction does not dominate all uses!
%1 = bitcast i8* %0 to %union._float3
%5 = extractvalue %union._float3 %1, 0
Instruction does not dominate all uses!
%5 = extractvalue %union._float3 %1, 0
%6 = extractvalue %struct.anon %5, 0
Instruction does not dominate all uses!
%6 = extractvalue %struct.anon %5, 0
%8 = fadd float %7, %6
Instruction does not dominate all uses!
%8 = fadd float %7, %6
store float %8, float addrspace(1)* %4, align 4
Broken module found, compilation aborted!

Yuri_K_Intel · ‎09-05-2013

Hi Chris, Indeed, I see this behavior on latest released version (XE 2013), but our internal development version builds this kernel just fine. So, please expect a fix to be available in next version of our SDK. I don't have any estimate about the date though. Thanks, Yuri

Yuri_K_Intel · ‎10-02-2013

Chris, Recently released version (XE 2013 R2, http://software.intel.com/en-us/vcsource/tools/opencl-sdk-xe) contains a fix for this issue. Could you please upgrade to it and verify the fix? Thanks, Yuri

Chris_Juelg · ‎10-02-2013

I just did the upgrade and the formerly problematic code runs like a charm.

Thank you for the info.