<div dir="ltr"><div><div><div><div>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:<br>char predicate = (op == JOIN_LT || op == JOIN_LE) ? oval < ival : oval > ival; <br>
</div><br>OpenCL platform:    Experiment Intel Gen OCL Driver<br>ASSERTION FAILED: Type is not supported by the instruction<br> 
 at file /root/WORK/beignet/backend/src/ir/context.cpp, function void 
gbe::ir::Context::append(const gbe::ir::Instruction&), line 160<br>Stack dump:<br>0.    Running pass 'Function Pass Manager' on module '/tmp/fileoWfSUJ.ll'.<br>1.    Running pass 'Gen Back-End' on function '@nlj_count_int'<br>
<br>gdb stack trace:<br>#0  gbe::onFailedAssertion (msg=<optimized out>, file=<optimized out>, <br>    fn=<optimized out>, line=<optimized out>)<br>    at /root/WORK/beignet/backend/src/sys/assert.cpp:76<br>
#1  0x00007f9d2b8090c2 in gbe::ir::Context::append (this=0x7f9ce4502f28, <br>    insn=...) at /root/WORK/beignet/backend/src/ir/context.cpp:160<br>#2  0x00007f9d2b85d0e9 in SEL<gbe::ir::Type, gbe::ir::Register, gbe::ir::Tuple><br>
    (this=0x7f9ce4502f28)<br>    at /root/WORK/beignet/backend/src/./ir/instruction.hxx:56<br>#3  SEL (src2=..., src1=..., src0=..., dst=..., type=gbe::ir::TYPE_BOOL, <br>    this=0x7f9ce4502f28) at /root/WORK/beignet/backend/src/./ir/context.hpp:150<br>
#4  gbe::GenWriter::emitSelectInst (this=0x7f9ce4502f00, I=...)<br>    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1871<br>#5  0x00007f9d2b86b795 in visitSelectInst (I=..., this=0x7f9ce4502f00)<br>    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:547<br>
#6  visitSelect (I=..., this=0x7f9ce4502f00)<br>    at /usr/local/include/llvm/IR/Instruction.def:165<br>#7  visit (I=..., this=<optimized out>)<br>    at /usr/local/include/llvm/IR/Instruction.def:165<br>#8  gbe::GenWriter::emitBasicBlock (this=this@entry=0x7f9ce4502f00, <br>
    BB=BB@entry=0x7f9ce40f68b0)<br>    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1020<br>#9  0x00007f9d2b86bba3 in gbe::GenWriter::emitFunction (<br>    this=this@entry=0x7f9ce4502f00, F=...)<br>    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1458<br>
#10 0x00007f9d2b8760d9 in runOnFunction (F=..., this=0x7f9ce4502f00)<br>    at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:487<br>#11 gbe::GenWriter::runOnFunction (this=0x7f9ce4502f00, F=...)<br>...<br>...<br>
<br></div>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.<br><br></div>Thanks,<br></div>/Ed<br><div>
<div><div><div><br><div><br><br>#define JOIN_LT      (-1)<br>#define JOIN_LE      (-2)<br>#define JOIN_EQ         0<br>#define JOIN_GT      1<br>#define JOIN_GE      2<br><br><br>__kernel void nlj_count_int(<br>    __global const int* outer,<br>
    const unsigned int outer_tuples,<br>    __global const int* inner,<br>    const unsigned int inner_tuples,<br>    __global unsigned int* temp_buffer,<br>    const unsigned int tuples_per_thread,<br>    const char op<br>
) {<br>    int pos = tuples_per_thread*get_global_id(0); <br>    unsigned int result_tuples = 0;    // Sum up the result tuples for this value.<br>    for (unsigned int i=0; i<tuples_per_thread; ++i) {<br>        if (pos >= outer_tuples)<br>
            break;<br>        // Check how many join partners this tuple has:<br>        int oval = outer[pos];<br>        for (unsigned int j=0; j<inner_tuples; ++j) {<br>            int ival = inner[j];<br>            // Evaluate the join predicate.<br>
            char predicate = (op == JOIN_LT || op == JOIN_LE) ? oval < ival : oval > ival; <br>            predicate |= (op == JOIN_LE || op == JOIN_GE) ? oval == ival : 0;<br>            predicate = (op == JOIN_EQ) ? oval == ival : predicate;<br>
            // If this matches, sum it up.<br>            result_tuples += predicate ? 1 : 0;    <br>        }<br>        // Move to the next tuple<br>        pos++;<br>    }<br>    temp_buffer[get_global_id(0)] = result_tuples; <br>
}<br><br></div></div></div></div></div></div>