[Beignet] Variety of assert failures with packed structs

[Lorenzo Pistone]

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

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20150924/550cafd5/attachment.html>