[Beignet] [BUG] piglit test case fail (get-global-size)
Boqun Feng
boqun.feng at intel.com
Sun Jun 23 21:14:51 PDT 2013
I found this bug is introduced by
commit 37687ac7c7b3b5bf5fb7ceeb9229c503afe74c87
the llvm IR for function get_global_size as follow:
get_global_size.exit2: ; preds = %7, %get_global_size.exit, %0
%9 = phi i1 [ true, %7 ], [ false, %get_global_size.exit ], [ false, %0 ]
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
this line cause the problem
%10 = phi i32 [ %5, %7 ], [ %5, %get_global_size.exit ], [ %1, %0 ]
%.0.i1 = phi i32 [ %8, %7 ], [ 1, %get_global_size.exit ], [ 1, %0 ]
%11 = mul i32 %.0.i1, %10
%12 = tail call ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
%13 = tail call ptx_device i32 @__gen_ocl_get_local_size0() nounwind readnone
%14 = tail call ptx_device i32 @__gen_ocl_get_group_id0() nounwind readnone
%15 = mul i32 %14, %13
%16 = add i32 %15, %12
br i1 %3, label %17, label %get_global_size.exit4
; <label>:17 ; preds = %get_global_size.exit2
%18 = tail call ptx_device i32 @__gen_ocl_get_global_size1() nounwind readnone
br label %get_global_size.exit4
get_global_size.exit4: ; preds = %17, %get_global_size.exit2
%.0.i3 = phi i32 [ %18, %17 ], [ 1, %get_global_size.exit2 ]
br i1 %9, label %19, label %get_global_size.exit6
; <label>:19 ; preds = %get_global_size.exit4
%20 = tail call ptx_device i32 @__gen_ocl_get_global_size2() nounwind readnone
br label %get_global_size.exit6
get_global_size.exit6: ; preds = %19, %get_global_size.exit4
%.0.i5 = phi i32 [ %20, %19 ], [ 1, %get_global_size.exit4 ]
%21 = mul i32 %.0.i3, %16
%22 = mul i32 %21, %.0.i5
br i1 %3, label %.thread17.i, label %get_global_id.exit
On Mon, Jun 24, 2013 at 11:34:22AM +0800, Boqun Feng wrote:
> 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;
> > }
>
More information about the Beignet
mailing list