AnsweredAssumed Answered

workgroups, scalar registers, constant buffers

Question asked by foomanchoo on Feb 6, 2013
Latest reply on Feb 7, 2013 by himanshu.gautam

good morning.


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[0];


uint coeff63 = coeff[63];


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.