[Beignet] Beignet code generator assertion on assignment with conditional operators

[Edward Ching]

I get the following error from Beignet when building the attached opencl
kernel with clBuildProgram(). Itooks like the problem is triggered by line
27 in the attached kernel:
char predicate = (op == JOIN_LT || op == JOIN_LE) ? oval < ival : oval >
ival;

OpenCL platform:    Experiment Intel Gen OCL Driver
ASSERTION FAILED: Type is not supported by the instruction
  at file /root/WORK/beignet/backend/src/ir/context.cpp, function void
gbe::ir::Context::append(const gbe::ir::Instruction&), line 160
Stack dump:
0.    Running pass 'Function Pass Manager' on module '/tmp/fileoWfSUJ.ll'.
1.    Running pass 'Gen Back-End' on function '@nlj_count_int'

gdb stack trace:
#0  gbe::onFailedAssertion (msg=<optimized out>, file=<optimized out>,
    fn=<optimized out>, line=<optimized out>)
    at /root/WORK/beignet/backend/src/sys/assert.cpp:76
#1  0x00007f9d2b8090c2 in gbe::ir::Context::append (this=0x7f9ce4502f28,
    insn=...) at /root/WORK/beignet/backend/src/ir/context.cpp:160
#2  0x00007f9d2b85d0e9 in SEL<gbe::ir::Type, gbe::ir::Register,
gbe::ir::Tuple>
    (this=0x7f9ce4502f28)
    at /root/WORK/beignet/backend/src/./ir/instruction.hxx:56
#3  SEL (src2=..., src1=..., src0=..., dst=..., type=gbe::ir::TYPE_BOOL,
    this=0x7f9ce4502f28) at
/root/WORK/beignet/backend/src/./ir/context.hpp:150
#4  gbe::GenWriter::emitSelectInst (this=0x7f9ce4502f00, I=...)
    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1871
#5  0x00007f9d2b86b795 in visitSelectInst (I=..., this=0x7f9ce4502f00)
    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:547
#6  visitSelect (I=..., this=0x7f9ce4502f00)
    at /usr/local/include/llvm/IR/Instruction.def:165
#7  visit (I=..., this=<optimized out>)
    at /usr/local/include/llvm/IR/Instruction.def:165
#8  gbe::GenWriter::emitBasicBlock (this=this at entry=0x7f9ce4502f00,
    BB=BB at entry=0x7f9ce40f68b0)
    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1020
#9  0x00007f9d2b86bba3 in gbe::GenWriter::emitFunction (
    this=this at entry=0x7f9ce4502f00, F=...)
    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1458
#10 0x00007f9d2b8760d9 in runOnFunction (F=..., this=0x7f9ce4502f00)
    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:487
#11 gbe::GenWriter::runOnFunction (this=0x7f9ce4502f00, F=...)
...
...

Is this a code gen bug or is there something illegal about that above line?
I'm using the Dec 11 repo with commit
625d5aa18446206cd6a00a52d8a2094948fd9d93 at the tip.

Thanks,
/Ed



#define JOIN_LT      (-1)
#define JOIN_LE      (-2)
#define JOIN_EQ         0
#define JOIN_GT      1
#define JOIN_GE      2


__kernel void nlj_count_int(
    __global const int* outer,
    const unsigned int outer_tuples,
    __global const int* inner,
    const unsigned int inner_tuples,
    __global unsigned int* temp_buffer,
    const unsigned int tuples_per_thread,
    const char op
) {
    int pos = tuples_per_thread*get_global_id(0);
    unsigned int result_tuples = 0;    // Sum up the result tuples for this
value.
    for (unsigned int i=0; i<tuples_per_thread; ++i) {
        if (pos >= outer_tuples)
            break;
        // Check how many join partners this tuple has:
        int oval = outer[pos];
        for (unsigned int j=0; j<inner_tuples; ++j) {
            int ival = inner[j];
            // Evaluate the join predicate.
            char predicate = (op == JOIN_LT || op == JOIN_LE) ? oval < ival
: oval > ival;
            predicate |= (op == JOIN_LE || op == JOIN_GE) ? oval == ival :
0;
            predicate = (op == JOIN_EQ) ? oval == ival : predicate;
            // If this matches, sum it up.
            result_tuples += predicate ? 1 : 0;
        }
        // Move to the next tuple
        pos++;
    }
    temp_buffer[get_global_id(0)] = result_tuples;
}
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20131219/a9ea087a/attachment.html>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: join.cl
Type: application/simple-filter+xml
Size: 1144 bytes
Desc: not available
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20131219/a9ea087a/attachment.bin>