clspv
clspv copied to clipboard
memcpy replacement unsupported case
Building the following kernel from SHOC Level0 DeviceMemory
{
int gid = get_global_id(0), num_thr = get_global_size(0), grpid=get_group_id(0), j = 0;
float sum = 0;
int tid=get_local_id(0), localSize=get_local_size(0), litems=4096/localSize, goffset=localSize*grpid+tid*litems;
int s = tid;
__local float lbuf[4096];
for ( ; j<litems && j<(size-goffset) ; ++j)
lbuf[tid*litems+j] = data[goffset+j];
for (int i=0 ; j<litems ; ++j,++i)
lbuf[tid*litems+j] = data[i];
barrier(CLK_LOCAL_MEM_FENCE);
for (j=0 ; j<3000 ; ++j) {
float a0 = lbuf[(s+0)&(4095)];
float a1 = lbuf[(s+1)&(4095)];
float a2 = lbuf[(s+2)&(4095)];
float a3 = lbuf[(s+3)&(4095)];
float a4 = lbuf[(s+4)&(4095)];
float a5 = lbuf[(s+5)&(4095)];
float a6 = lbuf[(s+6)&(4095)];
float a7 = lbuf[(s+7)&(4095)];
float a8 = lbuf[(s+8)&(4095)];
float a9 = lbuf[(s+9)&(4095)];
float a10 = lbuf[(s+10)&(4095)];
float a11 = lbuf[(s+11)&(4095)];
float a12 = lbuf[(s+12)&(4095)];
float a13 = lbuf[(s+13)&(4095)];
float a14 = lbuf[(s+14)&(4095)];
float a15 = lbuf[(s+15)&(4095)];
sum += a0+a1+a2+a3+a4+a5+a6+a7+a8+a9+a10+a11+a12+a13+a14+a15;
s = (s+16)&(4095);
}
output[gid] = sum;
}
with the following command line
clspv -w -descriptormap=descriptors.map source.cl -cluster-pod-kernel-args -cl-single-precision-constant -images=0 -pod-ubo -int8 -o compiled.spv
results in the following assertion being hit:
clspv: /path/to/clspv/lib/ReplaceLLVMIntrinsicsPass.cpp:214: bool {anonymous}::ReplaceLLVMIntrinsicsPass::replaceMemcpy(llvm::Module&): Assertion `isa<ConstantInt>(CI->getArgOperand(2))' failed.
Abandon