[Beignet] [BUG] kernel fails at the backend

[Zhigang Gong]

On Sun, Oct 06, 2013 at 01:30:15AM -0400, Matt Harvey wrote:
> Hi
> 
> I'm having trouble and was wondering if you guys had any insight. This
> simple kernel:
> __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
> CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

Try to change the above statement to:
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

> 
> __kernel void edge_detect(__read_only image2d_t image_in,
>                         __write_only image2d_t image_out)
> {
>     const int2 pos = { get_global_id(0), get_global_id(1) };
>     //make this a pass through
>     float4 data = read_imagef(image_in, sampler, pos);
>     write_imagef(image_out, pos, data);
> }
> 
> Fails after compilation with this error message:
> ASSERTION FAILED: arg != NULL
>   at file /media/matto/Programming/src/beignet/backend/src/ir/sampler.cpp,
> function void gbe::ir::SamplerSet::append(gbe::ir::Register,
> gbe::ir::Context*), line 60
> Stack dump:
> 0. Running pass 'Function Pass Manager' on module '/tmp/fileB926yJ.ll'.
> 1. Running pass 'Gen Back-End' on function '@edge_detect'

It seems that for the __constant/constant qualifier, the compiler
will not do constant propagation. The only worked qualifier for sampler_t
is const, if you want to define it in the kernel. We will fix this problem
in the futuer, before that, you can use const as a workaround.

Thanks for reporting this bug.

> 
> This looks like a bug to me.
> 
> Any ideas?
> 
> Thanks

> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet