4 Replies Latest reply on Feb 7, 2013 11:56 PM by himanshu.gautam

    workgroups, scalar registers, constant buffers

    foomanchoo

      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.

        • Re: workgroups, scalar registers, constant buffers
          coordz

          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.

           

          You appear to be using the global thread id in your calculation which varies per thread which leads to the values being put in VGPRs.