Alignment for LD/ST instructions may exceed HSAIL limit (BRIG_ALIGNMENT_256)

[enigmastudio]

I'm on the latest HLC compiler from the branch hsail-stable-3.7. I compiled with -O2.

The following kernel doesn't assemble after compilation, because the alignment info generated by the code generator can exceed HSAIL's limit of BRIG_ALIGNMENT_256:

__kernel void Bug(size_t index)
{
    __global uint *ptr = (__global uint *)(index*512);
    *ptr = 0xdeadc0de;
}

The emitted error is:

>   st_global_align(512)_u32    3735929054, [$d0];
>                   ^
input(16,18): Invalid alignment

ERROR:  The following command failed with return code 1.
        HSAILasm -o /tmp/hsa_finalizer-qVC9u8/temp.hsail /tmp/cloc31620/temp.hsail

I think all that needs to be done is to limit the alignment info produced by the code generator to BRIG_ALIGNMENT_256.