Use of "volatile" in warp safe programming


There is one thing I am unsure about.
I use the followig warp safe code (yes, I am warned about possible problems in the future).

volatile unsigned a[32];
unsigned *pD;   // points at global memory
if (threadIdx.x < 32) {
	if (threadIdx.x == 0) {
                 // some assignments, say:

		a[0] = 0;
		a[1] = 1;
		a[3] = 3;
		a[10] = 10;
		a[31] = 31;
	pD[threadIdx.x] = a[threadIdx.x];

Everything is inside one warp here. I know that to get correct behavior I must declare “a” as volatile to defend against compiler code rearrangements. But what about pD? Should it also be declared volatile?

Thank you.

Mike, Megaputer Intelligence Ltd.

a doesn’t need to be declared as volatile, the way you have shown it. a is a thread-local variable which means only one thread can access it. Therefore there is no risk of multi-thread access to a, and no particular need for volatile.

Mmm… There is multithread access. Thread 0 writes to “a” while all threads in this warp read it.

the correct verion is

shared volatile unsigned a[32];

“a” is shared array, not local.
Purifying the code from unnecessary details I threw off too much…

the pointer pD is a shared variable, too.

shared unsigned *pD;