# Can \_\_syncthreads exist in multiple control path ?

**URL:** <https://forums.developer.nvidia.com/t/can-syncthreads-exist-in-multiple-control-path/4696>\
**Category:** CUDA Programming and Performance\
**Created:** [July 23, 2008, 5:11am UTC](https://forums.developer.nvidia.com/t/can-syncthreads-exist-in-multiple-control-path/4696 "2008-07-23T05:11:49Z")\
**Posts on this page:** 4\
**Page:** 1

<div class="post-metadata">

**Author:** ![lee222](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@lee222](https://forums.developer.nvidia.com/u/lee222)\
**Post date:** [July 23, 2008, 5:11am UTC](https://forums.developer.nvidia.com/t/can-syncthreads-exist-in-multiple-control-path/4696/1 "2008-07-23T05:11:49Z")

</div>

Can \_\_syncthreads() exist in multiple control path?

CUDA Programming Guide book says “\_\_syncthread() is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects”.

For me, it seems that the last sentence gives some room for using \_\_syncthread() in multiple control paths.  
Suppose that execution path diverges like a tree, but it is merged back at some point later. If \_\_syncthreads() are distributed among the diverged paths, but no matter which path a thread traverse, the thread can see identical number of \_\_syncthreads() before reaching the merge point, then can this program run correctly?

---

<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:** [July 23, 2008, 3:44pm UTC](https://forums.developer.nvidia.com/t/can-syncthreads-exist-in-multiple-control-path/4696/2 "2008-07-23T15:44:56Z")

</div>

> [@](#):
>
> Can \_\_syncthreads() exist in multiple control path?
> 
> [snapback]414583[/snapback]

No. It’s not just the number of syncthreads, it’s which syncthreads as well.

---

<div class="post-metadata">

**Author:** ![lee222](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@lee222](https://forums.developer.nvidia.com/u/lee222)\
**Post date:** [July 24, 2008, 6:50am UTC](https://forums.developer.nvidia.com/t/can-syncthreads-exist-in-multiple-control-path/4696/3 "2008-07-24T06:50:03Z")

</div>

> [@](#):
>
> No. It’s not just the number of syncthreads, it’s which syncthreads as well.
> 
> [snapback]414842[/snapback]

Then, the following code will not work unless all threads in a block choose the same control path.

//tid is a threadID

```
    if( f(tid) ) { 

          __syncthreads();

    }

    else {

          __syncthreads();

    }

```

---

<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:** [July 24, 2008, 7:19am UTC](https://forums.developer.nvidia.com/t/can-syncthreads-exist-in-multiple-control-path/4696/4 "2008-07-24T07:19:36Z")

</div>

> [@](#):
>
> Then, the following code will not work unless all threads in a block choose the same control path.
> 
> //tid is a threadID
> 
> if( f(tid) ) {
> 
> \_\_syncthreads();
> 
> }
> 
> else {
> 
> \_\_syncthreads();
> 
> }
> 
> [snapback]415137[/snapback]

As far as I know, that is correct (that is, it won’t work).