When compiling a simple single function kernel (patch attached, also cut and pasted below), the compiler asserted on unsupported intrinsics:
ASSERTION FAILED: Unsupported intrinsics
at file
/root/WORK/test_split/beignet/backend/src/llvm/llvm_gen_backend.cpp,
function void gbe::GenWriter::regAllocateCallInst(llvm::CallInst&), line
1688
Stack dump:
0. Running pass 'Function Pass Manager' on module '/tmp/file0yciwa.ll'.
1. Running pass 'Gen Back-End' on function '@test_split'
It looks like this is triggered by the "shr_mem[pid+offset]=0;" line when 0
is assigned. Is this a bug? It looks like the kernel code is not doing
anything illegal. It compiled ok using Intel OpenCL SDK 2013 (on Linux, for
CPU)
Thanks,
/Ed
test_split.cl:
__kernel void
test_split(
__global int *in,
int stride,
int n_recs,
int n_parts,
__global int *out,
__local int *shr_mem)
{
int pid;
int glb_tid = get_global_id(0);
int loc_tid = get_local_id(0);
int offset = loc_tid*n_parts;
for(pid=0; pid<n_parts; pid++)
{
shr_mem[pid+offset]=0;
}
for(int pos=glb_tid; pos<n_recs; pos+=stride)
{
pid=in[pos];
shr_mem[pid+offset]++;
}
for(pid=0; pid<n_parts; pid++)
{
out[glb_tid + stride*pid] = shr_mem[pid+offset];
}
}
compiler_test_split.cpp:
void compiler_test_split(void)
{
OCL_CREATE_KERNEL("test_split");
}
0001-test-case-to-trigger-Unsupported-intrinsics-when-com.patch
Description: Binary data
_______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
