It seems that packed structs are completely broken in Beignet 1.1.0, Ivy
Bridge code gen. The following kernel
typedef struct __attribute__((packed)){
float v;
} point;
__kernel void test(__global point* points){
__global point* p = points + get_global_id(0);
p->v = 0.1f * p->v;
}
makes the compiler crash with a failed assertion:
ASSERTION FAILED: elemSize == GEN_BYTE_SCATTER_WORD || elemSize ==
GEN_BYTE_SCATTER_BYTE
at file
/home/pisto/sorgenti/Beignet-1.1.0-Source/backend/src/backend/gen_insn_selection.cpp,
function void
gbe::LoadInstructionPattern::emitUnalignedByteGather(gbe::Selection::Opaque&,
const gbe::ir::LoadInstruction&, uint32_t, gbe::GenRegister,
gbe::ir::BTI) const, line 3561
Replacing the type of point::v with float2 generates another assertion:
ASSERTION FAILED: count == 4 || count == 2
at file
/home/pisto/sorgenti/Beignet-1.1.0-Source/backend/src/./backend/gen_register.hpp,
function static gbe::GenRegister
gbe::GenRegister::splitReg(gbe::GenRegister, uint32_t, uint32_t),
line 284
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
