<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>