Global variables not being updated?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I’m writing an application which will use multiple thread blocks that will all access a global variable. However when this variable is changed, the changes seem to only be visible to the thread block which changed it. The variable labeled as device and I have seen this problem when I declare it as volatile as well:

__device__ volatile knode * tempKnodeG;

.

Additionally I have used a threadfence after changing the variable. A code snippet is below:

int thid = blockIdx.x * blockDim.x + threadIdx.x;

int bthid = threadIdx.x;

int i;

for(i = 0; i < height; i++){

	//do stuff...

	if(thread_should_update_value){

		//this will only occur for a single thread per each for loop iteration

		tempKnodeG = new_location;

		__threadfence();

	}

	//sync the thread blocks; wait is initialized to 0

	if(bthid == 0){

		atomicInc(&wait, UINT_MAX);

		volatile unsigned int * value = &wait;

		__threadfence();

		while(*value < NUM_BLOCKS*(i+1));

	}

	__syncthreads();

}

As you can see I’m syncing the thread blocks as well… This seems to be working fine. However once the for loop is done, tempKnodeG doesn’t seem to have the same value across thread blocks; whichever thread block that changed tempKnodeG sees the correct value, and the one(s) that didn’t don’t. It’s almost as if it is being treated as a shared variable.

Thanks for any help!