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.

Outcomes