[Beignet] [BUG] piglit test case fail (get-global-size)

[Boqun Feng]

Yes, I find the same problem, and I am trying to toggle on OUTPUT_LLVM
flag to see the compiler IR

On Mon, Jun 24, 2013 at 11:27:31AM +0800, Sun, Yi wrote:
> But it's a bit strange that I didn't found any bool variables in the
> test case after a quick looking through. The kernel is attached.
> 
> Thanks
>   --Sun, Yi
> 
> On Mon, 2013-06-24 at 11:20 +0800, Boqun Feng wrote:
> > This one is caused by bool variable's limited support of our compiler
> > 
> > On Mon, Jun 24, 2013 at 03:15:13AM +0000, Sun, Yi wrote:
> > > I'm not sure it is caused by the built-in function get-global-size
> > > specified. 
> > > Anyone has any idea?
> > > 
> > > reproduce step:
> > > [piglit]# bin/cl-program-tester
> > > tests/cl/program/execute/get-global-size.cl
> > > 
> > > 
> > > Output:
> > > ## Test: get_global_id
> > > (/GFX/Test/Piglit/piglit/tests/cl/program/program-tester.c) ##
> > > 
> > > # Platform supporting only version 1.0. Running test on that version.
> > > # Running on:
> > > #   Platform: Experiment Intel Gen OCL Driver
> > > #   Device: Intel HD Graphics Family
> > > #   OpenCL version: 1.1
> > > #   OpenCL C version: 1.0
> > > #   Build options: 
> > > ASSERTION FAILED: TODO Boolean values cannot escape their definition
> > > basic block
> > >   at
> > > file /home/works/opencl/beignet/backend/src/llvm/llvm_gen_backend.cpp,
> > > function void gbe::GenWriter::emitMovForPHI(llvm::BasicBlock*,
> > > llvm::BasicBlock*), line 844
> > > Stack dump:
> > > 0.      Running pass 'Function Pass Manager' on module
> > > '/tmp/file292g7Y.ll'.
> > > 1.      Running pass 'Gen Back-End' on function '@fill3d'
> > > 
> > > 
> > > Thanks
> > >   --Sun, Yi
> > > _______________________________________________
> > > Beignet mailing list
> > > Beignet at lists.freedesktop.org
> > > http://lists.freedesktop.org/mailman/listinfo/beignet
> 

> /*!
> [config]
> name: get_global_id
> clc_version_min: 10
> 
> [test]
> name: 1D, global_size 4 0 0
> kernel_name: fill1d
> dimensions: 1
> global_size: 4 0 0
> local_size: 1 0 0
> arg_out: 0 buffer int[4] 4 4 4 4
> 
> [test]
> name: 1D, global_size 4 0 0, local_size 2 0 0
> kernel_name: fill1d
> dimensions: 1
> global_size: 4 0 0
> local_size: 2 0 0
> arg_out: 0 buffer int[4] 4 4 4 4
> 
> [test]
> name: 1D, global_size 4 0 0, local_size 4 0 0
> kernel_name: fill1d
> dimensions: 1
> global_size: 4 0 0
> local_size: 4 0 0
> arg_out: 0 buffer int[4] 4 4 4 4
> 
> [test]
> name: 2D, global_size 4 4 0, local_size 1 1 0
> kernel_name: fill2d
> dimensions: 2
> global_size: 4 4 0
> local_size: 1 1 0
> arg_out: 0 buffer int[16] 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
> 
> [test]
> name: 2D, global_size 4 4 0, local_size 2 2 0
> kernel_name: fill2d
> dimensions: 2
> global_size: 4 4 0
> local_size: 2 2 0
> arg_out: 0 buffer int[16] 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
> 
> [test]
> name: 2D, global_size 4 4 0, local_size 4 4 0
> kernel_name: fill2d
> dimensions: 2
> global_size: 4 4 0
> local_size: 4 4 0
> arg_out: 0 buffer int[16] 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
> 
> [test]
> name: 3D, global_size 4 4 4, local_size 1 1 1
> kernel_name: fill3d
> dimensions: 3
> global_size: 4 4 4
> local_size: 1 1 1
> arg_out: 0 buffer int[64] 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64
> 
> [test]
> name: 3D, global_size 4 4 4, local_size 2 2 2
> kernel_name: fill3d
> dimensions: 3
> global_size: 4 4 4
> local_size: 2 2 2
> arg_out: 0 buffer int[64] 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64
> 
> [test]
> name: 3D, global_size 4 4 4, local_size 4 4 4
> kernel_name: fill3d
> dimensions: 3
> global_size: 4 4 4
> local_size: 4 4 4
> arg_out: 0 buffer int[64] 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 \
>                           64 64 64 64 64 64 64 64 64 64 64 64 64 64 64 64
> !*/
> 
> kernel void fill1d(global int* out) {
>     unsigned int size = get_global_size(0);
> 	unsigned int id = get_global_id(0);
>     out[id] = size;
> }
> 
> kernel void fill2d(global int* out) {
>     unsigned int size = get_global_size(0) * get_global_size(1);
>     unsigned int id = get_global_id(0)*get_global_size(1)+get_global_id(1);
> 	out[id] = size;
> }
> 
> kernel void fill3d(global int* out) {
>     unsigned int size = get_global_size(0) * get_global_size(1) * get_global_size(2);
> 	unsigned int id = (get_global_id(0) * (get_global_size(1) * get_global_size(2)))
>         + (get_global_id(1) * get_global_size(2)) + get_global_id(2);
>     out[id] = size;
> }