i currently use 64 constant uint in every thread of every workgroup in my kernel like so:
__constant uint * coeff = coeff_kernel_argument;
uint coeff1 = coeff;
uint coeff63 = coeff;
i am delighted to see that the 64 values end up in scalar registers (i am talking about Tahiti), as confirmed with the ISA dump.
now i would like to change it to:
uint * coeff = coeff_kernel_argument + get_global_id(0) / get_local_size(0) * 64;
unfortunately this results in the coeff?? variables to be stored in vgprs, even though their values are the same for each thread of a workgroup.
i am also beginning to rewrite my kernel in AMDIL for other reasons and as such a workaround that only works on the IL level would be suffice.
thanks a lot.