# order of execution in a divergent warp

**URL:** <https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435>\
**Category:** CUDA Programming and Performance\
**Created:** [January 7, 2012, 4:03pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435 "2012-01-07T16:03:12Z")\
**Posts on this page:** 20\
**Page:** 1

<div class="post-metadata">

**Author:** ![mikee111](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikee111](https://forums.developer.nvidia.com/u/mikee111)\
**Post date:** [January 7, 2012, 4:03pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/1 "2012-01-07T16:03:12Z")

</div>

I’ve encountered a serious issue and since all my approaches how to solve it failed, I have to ask for external assistance :) Basically, I have a warp where first thread first atomically subtracts a variable, saves the result to a shared memory and according to this result all other threads do some computation, the code looks pretty much like this:

```auto
__device__ void taskFinishSortSORT(const int& tid, const int& taskIdx, volatile Task* task, CUdeviceptr ppsBuf, const int start, const int end, volatile int& left, volatile int& right, const TaskType type)

{

	__shared__ volatile int sharedFinished[MaxBlockHeight]; // We need shared memory to distribute data from thread with tid == 0 to all the others in the same warp

	if(tid == 0)

	{

		int finished = atomicSub(&g_taskStack.tasks[taskIdx].unfinished, 1);

		sharedFinished[threadIdx.y] = finished;

	}

	if(sharedFinished[threadIdx.y] == 0) // Finish the whole sort

	{

		taskFinishSort(tid, taskIdx, task);

	}

}

```

What’s happening is that all threads from 1 to 31 will do the second if block, before the first (with tid == 0) is done and I have no idea how to avoid it. Of course if this happens the threads will choose to do taskFinishSort when finished is not actually 0 and all crashes and burns … taskFinishSort is a large function, I guess that’s important since when I replace taskFinishSort with some dummy function (with just a return or some few lines of code) it all works fine. Also, it works OK in debug build with debug info, it goes to hell in release. We also looked into the ptx, it seemed that it gets compiled in the correct order, also this decision which should run first should be done by the scheduler on the GPU, right?

So, what am I doing wrong? Thanks in advance for any advice :)

---

<div class="post-metadata">

**Author:** ![pasoleatis](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pasoleatis](https://forums.developer.nvidia.com/u/pasoleatis)\
**Post date:** [January 7, 2012, 4:22pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/2 "2012-01-07T16:22:34Z")

</div>

What about \_\_syncthreads();? It will make all the threads to wait until all the threads reach that point. Or maybe some threadfence functions (there is one for the shared memory) will put a barrier so that all threads have to wait until the memory writes are finished.

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 7, 2012, 4:32pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/3 "2012-01-07T16:32:51Z")

</div>

Is [font=“Courier New”]tid[/font] the same as [font=“Courier New”]threadIdx.x[/font], or is it [font=“Courier New”]threadIdx.y\*blockDim.x+threadIdx.x[/font]?

In the latter case it looks like your code wouldn’t work with more than one warp per block.

BTW. on compute capability 1.2 and higher you don’t need to go through shared memory:

```auto
__device__ void taskFinishSortSORT(const int& tid, const int& taskIdx, volatile Task* task, CUdeviceptr ppsBuf, const int start, const int end, volatile int& left, volatile int& right, const TaskType type)

{

        int finished=0;

if(threadIdx.x == 0)

                finished = (atomicSub(&g_taskStack.tasks[taskIdx].unfinished, 1) == 0);

if(__any(finished)) // Finish the whole sort

        {

                taskFinishSort(tid, taskIdx, task);

        }

}

```

---

<div class="post-metadata">

**Author:** ![mikee111](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikee111](https://forums.developer.nvidia.com/u/mikee111)\
**Post date:** [January 7, 2012, 9:59pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/4 "2012-01-07T21:59:10Z")

</div>

> [@](#):
>
> What about \_\_syncthreads();? It will make all the threads to wait until all the threads reach that point. Or maybe some threadfence functions (there is one for the shared memory) will put a barrier so that all threads have to wait until the memory writes are finished.

\_\_syncthreads will afaik sync warps in a block, not threads in a warp … and yes, tried it, did not work, there’s \_\_threadfence for making sure that your global and shared changes are visible to the device, this is not the problem

> [@](#):
>
> Is [font=“Courier New”]tid[/font] the same as [font=“Courier New”]threadIdx.x[/font], or is it [font=“Courier New”]threadIdx.y\*blockDim.x+threadIdx.x[/font]?
> 
> In the latter case it looks like your code wouldn’t work with more than one warp per block.
> 
> BTW. on compute capability 1.2 and higher you don’t need to go through shared memory:
> 
> ```auto
> __device__ void taskFinishSortSORT(const int& tid, const int& taskIdx, volatile Task* task, CUdeviceptr ppsBuf, const int start, const int end, volatile int& left, volatile int& right, const TaskType type)
> 
> {
> 
> int finished=0;
> 
> if(threadIdx.x == 0)
> 
> finished = (atomicSub(&g_taskStack.tasks[taskIdx].unfinished, 1) == 0);
> 
> if(__any(finished)) // Finish the whole sort
> 
> {
> 
> taskFinishSort(tid, taskIdx, task);
> 
> }
> 
> }
> 
> ```

it’s the same, we have persistent threads (and if a block is used it’s always [32,y]) … as for your example, i’ll try if it will make the compiler do the finished = (atomicSub(&g\_taskStack.tasks[taskIdx].unfinished) before doing if(\_\_any(finished)), since that’s what’s happening now and that is the real issue

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 7, 2012, 10:34pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/5 "2012-01-07T22:34:19Z")

</div>

Is [font=“Courier New”]taskFinishSortSORT()[/font] ever called within conditional code? Or from within a loop where the number of iterations differs between the threads? Do threads return early from the kernel or any of the (device) functions called from it?

---

<div class="post-metadata">

**Author:** ![mikee111](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikee111](https://forums.developer.nvidia.com/u/mikee111)\
**Post date:** [January 8, 2012, 10:45am UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/6 "2012-01-08T10:45:38Z")

</div>

> [@](#):
>
> Is [font=“Courier New”]taskFinishSortSORT()[/font] ever called within conditional code? Or from within a loop where the number of iterations differs between the threads? Do threads return early from the kernel or any of the (device) functions called from it?

no, no and … no … I know it’s not very precise, I’ve tried to put a printf at the start of the function and then in some other places, and it shows that at the start all threads are together and then 1…31 threads do the second if, then 0 does the first if and then the second …

---

<div class="post-metadata">

**Author:** ![mikee111](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikee111](https://forums.developer.nvidia.com/u/mikee111)\
**Post date:** [January 8, 2012, 11:29am UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/7 "2012-01-08T11:29:14Z")

</div>

well, apparently the problem was inlining, so putting **noinline** in front of taskFinishSort solved the problem … we’re not sure now if it’s all ok, but the simple test we had in place now works as expected

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 8, 2012, 12:41pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/8 "2012-01-08T12:41:07Z")

</div>

Seems like the compiler isn’t putting the reconvergence point into the optimal position, and deinlining is helping it to find that. We’ve already checked the most common causes for this though.

Which compute capability is your device? Is [font=“Courier New”]g\_taskStack[/font] in shared memory?

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 8, 2012, 12:52pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/9 "2012-01-08T12:52:31Z")

</div>

Are there any (perhaps hidden) [font=“Courier New”]return[/font] or [font=“Courier New”]exit[/font] statements in your code at all, even if never taken? That would influence the choice of reconvergence point as well.

---

<div class="post-metadata">

**Author:** ![mikee111](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikee111](https://forums.developer.nvidia.com/u/mikee111)\
**Post date:** [January 8, 2012, 2:13pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/10 "2012-01-08T14:13:44Z")

</div>

> [@](#):
>
> Seems like the compiler isn’t putting the reconvergence point into the optimal position, and deinlining is helping it to find that. We’ve already checked the most common causes for this though.
> 
> Which compute capability is your device? Is [font=“Courier New”]g\_taskStack[/font] in shared memory?

CC 2.0, g\_taskStack is in global

> [@](#):
>
> Are there any (perhaps hidden) [font=“Courier New”]return[/font] or [font=“Courier New”]exit[/font] statements in your code at all, even if never taken? That would influence the choice of reconvergence point as well.

well, we think there are no hidden returns, there’s definitely no exit … what exactly would you mean by hidden, like a code that will never execute (in some run of a program or never) but exists nonetheless?

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 8, 2012, 2:52pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/11 "2012-01-08T14:52:46Z")

</div>

When saying ‘hidden’ I was thinking of occurrences inside other functions, like assert() or so. However code that will never execute but exists nonetheless will be just as interesting, as that would influence the static analysis of the compiler just as any other code.

I’m running out of ideas though. As I assume you don’t want to post your complete code here, we might have to leave it as it is, if nobody else comes up with a suggestion.

It would be possible to confirm the placement of the reconvergence point in the disassembled cubin (with [font=“Courier New”]cuobjdump -sass[/font], then look for the operand of the SSY instruction), however that would tell nothing about why the compiler places it there.

---

<div class="post-metadata">

**Author:** ![mikee111](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikee111](https://forums.developer.nvidia.com/u/mikee111)\
**Post date:** [January 8, 2012, 3:02pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/12 "2012-01-08T15:02:07Z")

</div>

> [@](#):
>
> When saying ‘hidden’ I was thinking of occurrences inside other functions, like assert() or so. However code that will never execute but exists nonetheless will be just as interesting, as that would influence the static analysis of the compiler just as any other code.
> 
> I’m running out of ideas though. As I assume you don’t want to post your complete code here, we might have to leave it as it is, if nobody else comes up with a suggestion.
> 
> It would be possible to confirm the placement of the reconvergence point in the disassembled cubin (with [font=“Courier New”]cuobjdump -sass[/font], then look for the operand of the SSY instruction), however that would tell nothing about _why_ the compiler places it there.

thanks for all your pointers and/or advices, we will try to check for the reconvergence point in disassembled cubin, i’ve tried to compare ptx output for a variant with a large function call and a small function call, those were both the same, but since both included an actual function call (and not an inlined function) I guess ptx output does not include all optimizations … We tried that at the point where we did not know inlining was at fault

---

<div class="post-metadata">

**Author:** ![Sarnath](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/sarnath/32/9742_2.png) [@Sarnath](https://forums.developer.nvidia.com/u/Sarnath)\
**Post date:** [January 10, 2012, 10:02am UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/13 "2012-01-10T10:02:12Z")

</div>

The order of execution of sub-warps after a warp-divergence is UNDEFINED.  
Any assumptions on that order will result in future-incompatible-programs.

This is what NVIDIA has been telling for long. HTH

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 10, 2012, 12:04pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/14 "2012-01-10T12:04:39Z")

</div>

If this were true then it would be illegal to ever use \_\_syncthreads() again after a warp-divergence.

There have to be some reasonable guarantees about reconvergence, or the CUDA execution model wouldn’t work. Unfortunately the PTX ISA manual contradicts itself on this subject: Section 8.5 (both in version 2.3 and 3.0) says, just as you pointed out

> [@](#):
>
> Threads in a CTA execute together, at least in appearance, until they come to a conditional control construct such as a conditional branch, conditional function call, or conditional return. If threads execute down different control flow paths, the threads are called divergent.

Later, it states (emphasis mine)

> [@](#):
>
> For divergent control flow, the optimizing code generator automatically determines points of re-convergence. **Therefore, a compiler or code author targeting PTX can ignore the issue of divergent threads, but has the opportunity to improve performance** by marking branch points as uniform when the compiler or author can guarantee that the branch point is non-divergent.

The description of the bar.sync instruction of course clearly indicates that it behaves differently whether a warp is divergent or not.

So, to the letter, you are right. But the bold statement above (pun intended) would be wrong then.

The C Programming Guide on the other hand does not mention such a restriction of \_\_syncthreads(). Appendix B.6 states

> [@](#):
>
> \_\_syncthreads() 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.

It does not mention that use of \_\_syncthreads() is limited even after the conditional code section. So one would assume that Nvidia has made sure their algorithm to insert reconvergence points works well enough to always insert one after any conditional code construct emitted by the compiler. No?

Any clarification by Nvidia employees?

---

<div class="post-metadata">

**Author:** ![Sarnath](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/sarnath/32/9742_2.png) [@Sarnath](https://forums.developer.nvidia.com/u/Sarnath)\
**Post date:** [January 10, 2012, 2:54pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/15 "2012-01-10T14:54:39Z")

</div>

Tera,

I never said anything about re-convergence.  
I only said NVIDIA has stated earlier that “The **ORDER** of execution of sub-warps is UNDEFINED”.  
I was just providing this info as “FYI” – if that would be of some help.

HTH,  
Best Regards,  
Sarnath

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 15, 2012, 8:41am UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/16 "2012-01-15T08:41:26Z")

</div>

I hadn’t realized you were replying just to the thread title, not to the initial post.

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 15, 2012, 8:41am UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/17 "2012-01-15T08:41:26Z")

</div>

I hadn’t realized you were replying just to the thread title, not to the initial post.

---

<div class="post-metadata">

**Author:** ![Revan1](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Revan1](https://forums.developer.nvidia.com/u/Revan1)\
**Post date:** [January 15, 2012, 5:31pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/18 "2012-01-15T17:31:14Z")

</div>

What haven’t been said in the originaly post is that inside the function taskFinishSort() there is another single threaded code between if(tid == 0) {…}. We haven’t thought it important but as it shows out the compiler probably thinks it would be best to merge these two blocks.  
We have solved this issue with mikee111 by tricking the compiler into not optimizing the code. The gimmick is to use a different id in each single thread condition of the entire call stack. This way the compiler cannot merge the blocks and the code works as intended.

By the way if you take this problem to the limit there is no clean way to initialize shared memory with less than 32 thread. So I suppose that for the compiler the \_\_syncthreads() marks the point of reconvergence.

---

<div class="post-metadata">

**Author:** ![Revan1](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Revan1](https://forums.developer.nvidia.com/u/Revan1)\
**Post date:** [January 15, 2012, 5:31pm UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/19 "2012-01-15T17:31:14Z")

</div>

What haven’t been said in the originaly post is that inside the function taskFinishSort() there is another single threaded code between if(tid == 0) {…}. We haven’t thought it important but as it shows out the compiler probably thinks it would be best to merge these two blocks.  
We have solved this issue with mikee111 by tricking the compiler into not optimizing the code. The gimmick is to use a different id in each single thread condition of the entire call stack. This way the compiler cannot merge the blocks and the code works as intended.

By the way if you take this problem to the limit there is no clean way to initialize shared memory with less than 32 thread. So I suppose that for the compiler the \_\_syncthreads() marks the point of reconvergence.

---

<div class="post-metadata">

**Author:** ![Sarnath](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/sarnath/32/9742_2.png) [@Sarnath](https://forums.developer.nvidia.com/u/Sarnath)\
**Post date:** [January 16, 2012, 4:35am UTC](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435/20 "2012-01-16T04:35:19Z")

</div>

```auto
for(int i=threadIdx.x; i<SHARED_MEM_CAPACITY; i += blockDim.x)

{

  sharedArray[i] = INIT_VALUE;

}

__syncthreads();

```

Wont this initialize a shared array with less than 32 threads a block?

[Next page](https://forums.developer.nvidia.com/t/order-of-execution-in-a-divergent-warp/25435.md?page=2)