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@entry=0x7f9ce4502f00,
BB=BB@entry=0x7f9ce40f68b0)
at /root/WORK/beignet/backend/src/llvm/llvm_gen_backend.cpp:1020
#9 0x00007f9d2b86bba3 in gbe::GenWriter::emitFunction (
this=this@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;
}
join.cl
Description: application/simple-filter
_______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet
