[Beignet] Program compilation crash on Ivy Bridge

[Nicolas Bourdaud]

Hi!

Two weeks ago, I mentioned a problem of OpenCL program crashing during
compilation on IvyBridge while compiling fine on Haswell. I have been
advised to test using llvm/clang 3.5. So I have compiled beignet 1.0.0
using llvm/clang 3.5 but it turns out that I still get the same problem
on IvyBridge (Haswell still works):

> OpenCL platform name: Intel Gen OCL Driver
> OpenCL version: OpenCL 1.2 beignet 1.0.0
> OpenCL device: Intel(R) HD Graphics IvyBridge GT2
> Running OpenCL 1.2 version
> going to compile CL prog
> Call parameter type does not match function signature!
> i128 8404990
>  float  %8 = call i32 @_Z6islessff(float %7, i128 8404990) #3
> Both operands to a binary operator are not of the same type!
>   %18 = fmul ppc_fp128 0xM0000000000FF40060000000000000000, float %17
> Invalid insertelement operands!
>   %20 = insertelement <4 x float> %19, ppc_fp128 %18, i32 1
> Invalid insertelement operands!
>   %26 = insertelement <4 x float> %25, ppc_fp128 0xM00000000000000000000000000000000, i32 3
> LLVM ERROR: Broken function found, compilation aborted!


I don't know what I could check. So any lead is welcome. For reference
the offending CPUs are:
 - Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz
 - Intel(R) Core(TM) i5-3470 CPU @ 3.20GHz
I have put the program provoking the crash in attachment (test.cl)

Thanks in advance to anyone who could shed some light on these issues.

Cheers,

Nicolas


On 11/18/2014 10:00 AM, Nicolas Bourdaud wrote:
> Thanks! I will try this
> 
> Nicolas
> 
> On 18/11/2014 02:52, Zhigang Gong wrote:
>> The following error log is from LLVM internal. It's likely a llvm/clang 3.4's bug.
>> The kernel works fine with both beignet 0.9.3 and 1.0.0 by using llvm/clang 3.5.
>> The recommended llvm version for beignet is LLVM 3.5.
>>
>> OpenCL platform name: Intel Gen OCL Driver
>> OpenCL version: OpenCL 1.2 beignet 0.9.3
>> OpenCL device: Intel(R) HD Graphics IvyBridge GT2
>> Running OpenCL 1.2 version
>> going to compile CL prog
>> Call parameter type does not match function signature!
>> i128 8404990
>>  float  %8 = call i32 @_Z6islessff(float %7, i128 8404990) #7
>> Both operands to a binary operator are not of the same type!
>>   %18 = fmul ppc_fp128 0xM0000000000FF40060000000000000000, float %17
>> Invalid insertelement operands!
>>   %20 = insertelement <4 x float> %19, ppc_fp128 %18, i32 1
>> Invalid insertelement operands!
>>   %26 = insertelement <4 x float> %25, ppc_fp128 0xM00000000000000000000000000000000, i32 3
>> Broken module found, compilation aborted!
>> Aborted
>>
>>> -----Original Message-----
>>> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
>>> Nicolas Bourdaud
>>> Sent: Monday, November 17, 2014 7:27 PM
>>> To: beignet at lists.freedesktop.org
>>> Subject: [Beignet] Program compilation crash on Ivy Bridge
>>>
>>> Hi!
>>>
>>> In beignet 0.9.3, beignet aborts in the middle on compilation with compiling for
>>> ivybridge microarchitecture, while it compiles and runs fine on Haswell.
>>>
>>> I have provided a test program (test.cl) that shows the problem and the
>>> generate log on haswell (haswell.log) and ivybridge (ivybridge.log). I have
>>> tested on different ivybridge CPU  (HD4000, HD2500) and the issue is the
>>> same. I have also tested on different haswell CPU (only HD4600) and it always
>>> goes fine.
>>>
>>> For information beignet has been tested for:
>>>  - ivybridge using linux 3.14 compiled with llvm 3.4.
>>>  - haswell using linux 3.14 compiled with llvm 3.4.
>>>  - haswell using linux 3.16 compiled with llvm 3.5 In all case, the program
>>> compiled and runs correctly with haswell, not with ivybridge.
>>>
>>> I could give a try with beignet 1.0.0, but maybe you the problem would be
>>> obvious to you.
>>>
>>> Any idea about the problem?
>>>
>>> Cheers,
>>>
>>> Nicolas Bourdaud
>>
> 
> 

-------------- next part --------------
// vim: set syntax=opencl:

#define IMSTRIDE	47

inline static
float4 transform_pixel(float4 a4)
{
	float4 r4;
	int is_null;
	float C, v;

	C = a4.x - a4.y;

	is_null = isless(C, 0.5f);


	v = select(a4.z, 0.0f, is_null);

	r4.x = a4.x;
	r4.y = 255.0f*v;
	r4.z = a4.y;
	r4.w = 0.0f;

	return r4;
}


__kernel void transform(__global const uchar4* restrict in,
                        __global uchar4* restrict out)
{
	int gix = get_global_id(0);
	int giy = get_global_id(1);

	float4 pix;
	uchar4 pix_in;

	pix_in = in[gix + giy*(IMSTRIDE/sizeof(*in))];
	pix = convert_float4(pix_in);
	pix = transform_pixel(pix);
	out[gix + giy*(IMSTRIDE/sizeof(*out))] = convert_uchar4_sat(pix);
}