# Global thread barrier

**URL:** <https://forums.developer.nvidia.com/t/global-thread-barrier/8498>\
**Category:** CUDA Programming and Performance\
**Created:** [March 23, 2009, 5:40pm UTC](https://forums.developer.nvidia.com/t/global-thread-barrier/8498 "2009-03-23T17:40:26Z")\
**Posts on this page:** 1\
**Showing post:** 58

<div class="post-metadata">

**Author:** ![Lev](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Lev](https://forums.developer.nvidia.com/u/Lev)\
**Post date:** [April 15, 2010, 4:25pm UTC](https://forums.developer.nvidia.com/t/global-thread-barrier/8498/58 "2010-04-15T16:25:29Z")

</div>

peastman, andradx

I think I found good and fast solution of your iteration problem.

**global** blocknum=0;  
**global** completedblockquantity=0;

at the start of the block

**shared** current\_block\_num;

**shared** int currentiteration, blockindex;

if (threadIdx==0)  
{  
current\_block\_num=atomic\_inc(blocknum);

currentiteration=current\_block\_num/blocks\_per\_iteration;  
blockindex=current\_block\_num%blocks\_per\_iteration;

some code for waiting prior blocks to complete.  
like this

wait while (completedblockquantity\>=currentiteration\*blocks\_per\_iteration)  
}

\_\_syncthreads();

data=datarray[blockindex\*blocksize+threadIdx.x]

and so on

at the end of the block

\_\_threadfence();

atomic\_inc(completedblockquanity);

at kernell launch you need to mass blocks according to (data size)\*(iteration quantity)

---

_[View the full topic](https://forums.developer.nvidia.com/t/global-thread-barrier/8498)._