cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

boxerab
Challenger

Question about global pointers

The AMD optimization guide states:

--------------------------------------------------------------------------------------------------------

Avoid writing code with dynamic pointer assignment on the GPU. For

example:

kernel void dyn_assign(global int* a, global int* b, global int* c)

{

global int* d;

size_t idx = get_global_id(0);

if (idx & 1) {

d = b;

} else {

d = c;

}

a[idx] = d[idx];

}

------------------------------------------------------------------------------------------------------

Now, in my case, I have the following code:

void KERNEL run(GLOBAL const uint* restrict idata, GLOBAL uchar* restrict odata) {

   ....

    idata +=  get_local_id(0);

    odata += get_local_id(0)

   .....

}

Is this a case of dynamic pointer assignment, since I am modifying the global pointers idata and odata?

i.e. should I use an index into idata, instead of modifying it?

Thanks.

0 Likes
3 Replies
dipak
Big Boss

Hi,

If you read that section carefully, it says that compiler must be aware of the base pointer of the data buffer (see bold line below) in order to perform certain optimizations.


This is inefficient because the GPU compiler must know the base pointer that every load comes from and in this situation, the compiler cannot determine what ‘d’ points to. So, both B and C are assigned to the same GPU resource, removing the ability to do certain optimizations.


I guess, in your case, it is not the same as above.

Regards,

0 Likes

Hmmm, I don't see much of a difference though. Both idata & odata are unknown to the compiler. Furthermore, idata is declared constant pointer and then changed :-(. Referencing through indexes, would at least cure constant aspect.

PS: Haven't read optimization guide, yet.

0 Likes

Thanks, guys. @nibal: the const modifier in this case indicates that what

is pointed to is constant, i.e. read only. The pointer itself can change.

The following incantation:

global uchar* const idata

would indicate that the pointer itself is constant.

Aaron

0 Likes