# Coalescing question possibly a write commit question

**URL:** <https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443>\
**Category:** CUDA Programming and Performance\
**Created:** [March 20, 2009, 7:41pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443 "2009-03-20T19:41:56Z")\
**Posts on this page:** 10\
**Page:** 1

<div class="post-metadata">

**Author:** ![JohnW](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JohnW](https://forums.developer.nvidia.com/u/JohnW)\
**Post date:** [March 20, 2009, 7:41pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/1 "2009-03-20T19:41:56Z")

</div>

I am getting some strange behavior when trying to read global memory. My application performs operations on elements in a large array, which is initialized to zero. The device code sets the first element of the data array to a value, then all subsequent data elements are copied from that element within the data elements array. I don’t expect all the values to be set, there are multiple cores running simultaneously so until the first block of threads complete the initial element will still be zero. The problem is that it’s accuracy is not that great, with a different number of errors each run.

[codebox]#include \<stdio.h\>

#include \<cuda\_runtime.h\>

#define mesh 512

const dim3 dimBlock(512);

const dim3 dimGrid(mesh/dimBlock.x, mesh/dimBlock.y, 1);

const int Num\_Elements = mesh\*mesh;

static long \*data\_elements\_h, \*data\_elements\_d = NULL;

// Device Function

**global** void Process\_Elements(long\* data\_elements\_d)

{

```
int ThreadID = (blockIdx.y*gridDim.x+blockIdx.x)*blockDim.y*blockDim.x+thre

```

adIdx.y\*blockDim.x+threadIdx.x;

```
if (ThreadID==0) data_elements_d[ThreadID] = 0xF;

else data_elements_d[ThreadID] = data_elements_d[0]; // Suspect read/write

__syncthreads();

```

}

// Host Function

#if **DEVICE\_EMULATION**

bool InitCudaDevice(void) { return true; }

#else

bool InitCudaDevice(void) { cudaSetDevice(0); return true; }

#endif

int main(int argc, char\* argv)

{

```
long input_h = 0xF;

int elements_error = 0;

if (!InitCudaDevice()) return 0;

size_t Data_size = Num_Elements*sizeof(long);

data_elements_h = (long *)malloc(Data_size);

data_elements_h[0] = input_h; // Set first element to the input value

for (int ThreadID=1; ThreadID<Num_Elements; ThreadID++) data_elements_h[ThreadID] = 0; // Init other elements to zero

cudaMalloc((void **) &data_elements_d, Data_size);

cudaMemcpy(data_elements_d, data_elements_h, Data_size, cudaMemcpyHostToDevice);

Process_Elements<<<dimGrid, dimBlock, 0, 0>>>(data_elements_d);

cudaMemcpy(data_elements_h, data_elements_d, Num_Elements*sizeof(long), cudaMemcpyDeviceToHost);

for (int k=0; k<Num_Elements; k++) {

	if (data_elements_h[k] != input_h) {

		elements_error++;

		//printf("data_elements_h(%d): %llX\n", k, data_elements_h[k]);

	}

}

if (elements_error) printf("%d OK, %d Errors out of %d\n", Num_Elements-elements_error, elements_error, Num_Elements); else printf("No Element Errors\n");

free(data_elements_h);

cudaFree(data_elements_d);

cudaThreadExit();

return 0;

```

}[/codebox]

If I initialize the data with the first element already set then I get no errors and every data element gets the value 0xF after one iteration. However I need the code to run on the device without having to copy all the memory back and forth.

How do I determine how what span to leave before reading a value written by a previous thread? I’d like to be able to reference data based off the ThreadID like so:

data\_elements\_d[ThreadID] = data\_elements\_d[ThreadID - padding];

Thanks for any pointers,

John

---

<div class="post-metadata">

**Author:** ![tmurray](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tmurray](https://forums.developer.nvidia.com/u/tmurray)\
**Post date:** [March 20, 2009, 8:04pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/2 "2009-03-20T20:04:21Z")

</div>

Check out \_\_threadfence() in the 2.2 beta.

---

<div class="post-metadata">

**Author:** ![JohnW](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JohnW](https://forums.developer.nvidia.com/u/JohnW)\
**Post date:** [March 20, 2009, 8:53pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/3 "2009-03-20T20:53:22Z")

</div>

Thanks! That was easier than expected. I’ve just registered for the Beta.

Just to check my understanding. The following should allow me to access data written by a previous block:

[codebox] **global** void Process\_Elements(long\* data\_elements\_d){

```
int ThreadID = (blockIdx.y*gridDim.x+blockIdx.x)*blockDim.y*blockDim.x+thre

```

adIdx.y\*blockDim.x+threadIdx.x;

```
if (ThreadID<512) data_elements_d[ThreadID] = 0xF;

else data_elements_d[ThreadID] = data_elements_d[ThreadID-512]; // Suspect read/write

__threadfence_block();

__syncthreads();

```

}[/codebox]

But this WOULDN’T allow me to access data by the previous thread (but would be much slower anyway):

[codebox] **global** void Process\_Elements(long\* data\_elements\_d){

```
int ThreadID = (blockIdx.y*gridDim.x+blockIdx.x)*blockDim.y*blockDim.x+thre

```

adIdx.y\*blockDim.x+threadIdx.x;

```
if (ThreadID==0) data_elements_d[ThreadID] = 0xF;

else data_elements_d[ThreadID] = data_elements_d[ThreadID-1]; // Suspect read/write

__threadfence();

__syncthreads();

```

}[/codebox]

Thanks again!

---

<div class="post-metadata">

**Author:** ![JohnW](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JohnW](https://forums.developer.nvidia.com/u/JohnW)\
**Post date:** [March 25, 2009, 4:26am UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/4 "2009-03-25T04:26:39Z")

</div>

Just to clarify, my above post was a question based on reading the (much improved and way better) Version 2.2 Beta Programming guide. I don’t have the 2.2 Beta myself, so don’t interpret the above code as correct.

If anyone does, I’d be most grateful for any confirmation.

Cheers,  
John

---

<div class="post-metadata">

**Author:** ![MisterAnderson42](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MisterAnderson42](https://forums.developer.nvidia.com/u/MisterAnderson42)\
**Post date:** [March 25, 2009, 2:09pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/5 "2009-03-25T14:09:34Z")

</div>

> [@](#):
>
> How do I determine how what span to leave before reading a value written by a previous thread? I’d like to be able to reference data based off the ThreadID like so:

How do you define previous? In CUDA, threads are run in a completely non-deterministic order.

---

<div class="post-metadata">

**Author:** ![JohnW](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JohnW](https://forums.developer.nvidia.com/u/JohnW)\
**Post date:** [March 25, 2009, 6:52pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/6 "2009-03-25T18:52:11Z")

</div>

> [@](#):
>
> How do you define previous? In CUDA, threads are run in a completely non-deterministic order.

With a large data set, larger than the number of threads times the number of MP (512\*30) the execution is divided into blocks. The “first” block is for threadid 0 up to blockIdx.x \* blockDim.x + threadIdx.x with the blockId and blockDim starting at zero and incrementing untill the entire data set is calculated.

My question is wether data calculated and updated in block 0 is available to “later” blocks, block 1000 for example, and if so how many blocks should I leave before assuming the data has been updated.

---

<div class="post-metadata">

**Author:** ![MisterAnderson42](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MisterAnderson42](https://forums.developer.nvidia.com/u/MisterAnderson42)\
**Post date:** [March 25, 2009, 8:14pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/7 "2009-03-25T20:14:26Z")

</div>

> [@](#):
>
> My question is wether data calculated and updated in block 0 is available to “later” blocks, block 1000 for example, and if so how many blocks should I leave before assuming the data has been updated.

You cannot make any assumptions on the order in which blocks are executed. The hardware is free to run block 1000 first if it chooses to do so.

---

<div class="post-metadata">

**Author:** ![E.D\_Riedijk](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@E.D\_Riedijk](https://forums.developer.nvidia.com/u/E.D_Riedijk)\
**Post date:** [March 25, 2009, 9:07pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/8 "2009-03-25T21:07:38Z")

</div>

> [@](#):
>
> You cannot make any assumptions on the order in which blocks are executed. The hardware is free to run block 1000 first if it chooses to do so.

in fact, i currently have a kernel that uses atomicInc to update the index where I write in my output array (I write from very few threads, compared to the number of threads running).

When looking at the output array, it seems that ‘higher number’ threads are running before ‘lower number’ threads (threadIdx.x + blockIdx.x\*blockDim.x).

---

<div class="post-metadata">

**Author:** ![JohnW](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JohnW](https://forums.developer.nvidia.com/u/JohnW)\
**Post date:** [March 25, 2009, 9:41pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/9 "2009-03-25T21:41:28Z")

</div>

> [@](#):
>
> You cannot make any assumptions on the order in which blocks are executed. The hardware is free to run block 1000 first if it chooses to do so.

Understood, thanks for the clarification. Will this change in 2.2 with \_\_threadfence();?

---

<div class="post-metadata">

**Author:** ![tmurray](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tmurray](https://forums.developer.nvidia.com/u/tmurray)\
**Post date:** [March 25, 2009, 10:18pm UTC](https://forums.developer.nvidia.com/t/coalescing-question-possibly-a-write-commit-question/8443/10 "2009-03-25T22:18:50Z")

</div>

Nope.