}
void GenContext::emitIndirectMoveInstruction(const SelectionInstruction &insn) {
- const GenRegister src = GenRegister::unpacked_uw(ra->genReg(insn.src(0)).nr, 0);
+ GenRegister src = ra->genReg(insn.src(0));
+ if(isScalarReg(src.reg()))
+ src = GenRegister::retype(src, GEN_TYPE_UW);
+ else
+ src = GenRegister::unpacked_uw(src.nr, src.subnr / typeSize(GEN_TYPE_UW));
+
const GenRegister dst = ra->genReg(insn.dst(0));
const GenRegister a0 = GenRegister::addr8(0);
uint32_t simdWidth = p->curr.execWidth;
compiler_function_constant0(__constant short *c0, __constant char *c1, __global int *dst, int value)
{
int id = (int)get_global_id(0);
- dst[id] = value + c0[id%69] + c1[15];
+ dst[id] = value + c0[id%69] + c1[0];
}
// Setup kernel and buffers
OCL_CREATE_KERNEL("compiler_function_constant0");
OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), NULL);
- OCL_CREATE_BUFFER(buf[1], 0, 256 * sizeof(char), NULL);
+ OCL_CREATE_BUFFER(buf[1], 0, 1 * sizeof(char), NULL);
OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(uint32_t), NULL);
OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
OCL_UNMAP_BUFFER(0);
OCL_MAP_BUFFER(1);
- for(uint32_t i = 0; i < 256; ++i)
- ((char *)buf_data[1])[i] = 10;
- ((char *)buf_data[1])[15] = 15;
+ ((char *)buf_data[1])[0] = 15;
OCL_UNMAP_BUFFER(1);
// Run the kernel