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?
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.