Re: [og7] Update nvptx_fork/join barrier placement

[Cesar Philippidis <cesar at codesourcery dot com>]

On 03/19/2018 07:04 AM, Tom de Vries wrote:
> On 03/09/2018 05:55 PM, Cesar Philippidis wrote:
>> On 03/09/2018 08:21 AM, Tom de Vries wrote:
>>> On 03/09/2018 12:31 AM, Cesar Philippidis wrote:
>>>> Nvidia Volta GPUs now support warp-level synchronization.
>>>
>>> Well, let's try to make that statement a bit more precise.
>>>
>>> All Nvidia architectures have supported synchronization of threads in a
>>> warp on a very basic level: by means of convergence (and unfortunately,
>>> we've seen that this is very error-prone).
>>>
>>> What is new in ptx 6.0 combined with sm_70 is the ability to sync
>>> divergent threads without having to converge, f.i. by using new
>>> instructions bar.warp.sync and barrier.sync.
>>
>> Yes. The major difference sm_70 GPU architectures and earlier GPUs is
>> that sm_70 allows the user to explicitly synchronize divergent warps. At
>> least on Maxwell and Pascal, the PTX SASS compiler uses two instructions
>> to branch, SYNC and BRA. I think, SYNC guarantees that a warp is
>> convergent at the SYNC point, whereas BRA makes no such guarantees.
>>
> 
> If you want to understand the interplay of sync (or .s suffix), branch
> and ssy, please read
> https://people.engr.ncsu.edu/hzhou/ispass_15-poster.pdf .

Interesting, thanks!

>> What's worse, once a warp has become divergent on sm_60 and earlier
>> GPUs, there's no way to reliably reconverge them. So, to avoid that
>> problem, it critical that the PTX SASS compiler use SYNC instructions
>> when possible. Fortunately, bar.warp.sync resolves the divergent warp
>> problem on sm_70+.
>>
>>>> As such, the
>>>> semantics of legacy bar.sync instructions have slightly changed on
>>>> newer
>>>> GPUs.
>>>
>>> Before in ptx 3.1, we have for bar.sync:
>>> ...
>>> Barriers are executed on a per-warp basis as if all the threads in a
>>> warp are active. Thus, if any thread in a warp executes a bar
>>> instruction, it is as if all the threads in the warp have executed
>>> the bar instruction. All threads in the warp are stalled until the
>>> barrier completes, and the arrival count for the barrier is incremented
>>> by the warp size (not the number of active threads in the warp). In
>>> conditionally executed code, a bar instruction should only be used if it
>>> is known that all threads evaluate the condition identically (the warp
>>> does not diverge).
>>> ...
>>>
>>> But in ptx 6.0, we have:
>>> ...
>>> bar.sync is equivalent to barrier.sync.aligned
>>> ...
>>> and:
>>> ...
>>> Instruction barrier has optional .aligned modifier. When specified, it
>>> indicates that all threads in CTA will execute the same barrier
>>> instruction. In conditionally executed code, an aligned barrier
>>> instruction should only be used if it is known that all threads in
>>> CTA evaluate the condition identically, otherwise behavior is undefined.
>>> ...
>>>
>>> So, in ptx 3.1 bar.sync should be executed in convergent mode (all the
>>> threads in each warp executing the same). But in ptx 6.0, bar.sync
>>> should be executed in the mode that the whole CTA is executing the same
>>> code.
>>>
>>> So going from the description of ptx, it seems indeed that the semantics
>>> of bar.sync has changed. That is however surprising, since it would
>>> break the forward compatibility that AFAIU is the idea behind ptx.
>>>
>>> So for now my hope is that this is a documentation error.
>>
>> I spent a lot of time debugging deadlocks with the vector length changes
>> and I have see no changes in the SASS code generated in the newer Nvidia
>> drivers when compared to the older ones, at lease with respect to the
>> barrier instructions. This isn't the first time I've seen
>> inconsistencies with thread synchronization in Nvidia's documentation.
>> For the longest time, the "CUDA Programming Guide" provided slightly
>> conflicting semantics for the __syncthreads() function, which ultimately
>> gets implemented as bar.sync in PTX.
>>
>>>> The PTX JIT will now, occasionally, emit a warpsync instruction
>>>> immediately before a bar.sync for Volta GPUs. That implies that warps
>>>> must be convergent on entry to those threads barriers.
>>>>
>>>
>>> That warps must be convergent on entry to bar.sync is already required
>>> by ptx 3.1.
>>>
>>> [ And bar.warp.sync does not force convergence, so if the warpsync
>>> instruction you mention is equivalent to bar.warp.sync then your
>>> reasoning is incorrect. ]
>>
>> I'm under the impression that bar.warp.sync converges all of the
>> non-exited threads in a warp.
> 
> I have not played around with the instruction yet, so I'm not sure, but
> what I read from the docs is that bar.warp.sync converges all of the
> non-exited threads in a warp only and only if it's positioned at a point
> post-dominating a divergent branch.
> 
> Consider this case:
> ...
> if (tid.x == 0)
>   {
>     A;
>     bar.warp.sync 32;
>     B;
>   }
> else
>   {
>     C;
>     bar.warp.sync 32;
>     D;
>   }
> ...
> AFAIU, this allows bar.warp.sync to synchronize the threads in the warp,
> _without_ converging.

I think that's partially wrong. Check out the literature for CUDA 9
cooperative groups, such as
<https://devblogs.nvidia.com/cooperative-groups/>, to get an idea of the
intent behind bar.warp.sync.

>> You'd still need to use bar.sync or some
>> variant of the new barrier instruction to converge the entire CTA. But
>> at the moment, we're still generating code that's backwards compatible
>> with sm_30.
>>
>>>> The problem in og7, and trunk, is that GCC emits barrier
>>>> instructions at
>>>> the wrong spots. E.g., consider the following OpenACC parallel region:
>>>>
>>>>     #pragma acc parallel loop worker
>>>>     for (i = 0; i < 10; i++)
>>>>       a[i] = i;
>>>>
>>>> At -O2, GCC generates the following PTX code:
>>>>
>>>>           {
>>>>                   .reg.u32        %y;
>>>>                   mov.u32 %y, %tid.y;
>>>>                   setp.ne.u32     %r76, %y, 0;
>>>>           }
>>>>           {
>>>>                   .reg.u32        %x;
>>>>                   mov.u32 %x, %tid.x;
>>>>                   setp.ne.u32     %r75, %x, 0;
>>>>           }
>>>>           @%r76   bra.uni $L6;
>>>>           @%r75   bra     $L7;
>>>>                   mov.u64 %r67, %ar0;
>>>>           // fork 2;
>>>>                   cvta.shared.u64 %r74, __oacc_bcast;
>>>>                   st.u64  [%r74], %r67;
>>>> $L7:
>>>> $L6:
>>>>           @%r75   bra     $L5;
>>>>           // forked 2;
>>>>                   bar.sync        0;
>>>>                   cvta.shared.u64 %r73, __oacc_bcast;
>>>>                   ld.u64  %r67, [%r73];
>>>>                   mov.u32 %r62, %ntid.y;
>>>>                   mov.u32 %r63, %tid.y;
>>>>                   setp.gt.s32     %r68, %r63, 9;
>>>>           @%r68   bra     $L2;
>>>>                   mov.u32 %r55, %r63;
>>>>                   cvt.s64.s32     %r69, %r62;
>>>>                   shl.b64 %r59, %r69, 2;
>>>>                   cvt.s64.s32     %r70, %r55;
>>>>                   shl.b64 %r71, %r70, 2;
>>>>                   add.u64 %r58, %r67, %r71;
>>>> $L3:
>>>>                   st.u32  [%r58], %r55;
>>>>                   add.u32 %r55, %r55, %r62;
>>>>                   add.u64 %r58, %r58, %r59;
>>>>                   setp.le.s32     %r72, %r55, 9;
>>>>           @%r72   bra     $L3;
>>>> $L2:
>>>>                   bar.sync        1;
>>>>           // joining 2;
>>>> $L5:
>>>>           // join 2;
>>>>           ret;
>>>>
>>>> Note the bar.sync instructions placed immediately after the forked
>>>> comment and before the joining comment. The problem here is that branch
>>>> above the forked comment guarantees that the warps are not synchronous
>>>> (when vector_length > 1, which is always the case).
>>>
>>> This is already advised against in ptx 3.1, so yes, we should fix this.
>>>
>>>> Likewise, bar.sync
>>>> instruction before joining should be placed after label L5 in order to
>>>> allow all of the threads in the warp to reach it.
>>>>
>>>
>>> Agreed.
>>>
>>>> The attached patch teaches the nvptx to make those adjustments.
>>>
>>> Can you show me a diff of the ptx for the test-case above for trunk?
>>
>> --- w-old.s     2018-03-08 15:19:47.139516578 -0800
>> +++ w.s 2018-03-09 08:42:52.217057332 -0800
>> @@ -46,9 +46,9 @@
>>                  st.u64  [%r74], %r67;
>>   $L7:
>>   $L6:
>> -       @%r75   bra     $L5;
>>          // forked 2;
>>                  bar.sync        0;
>> +       @%r75   bra     $L5;
>>                  cvta.shared.u64 %r73, __oacc_bcast;
>>                  ld.u64  %r67, [%r73];
>>                  mov.u32 %r62, %ntid.y;
>> @@ -68,9 +68,9 @@
>>                  setp.le.s32     %r72, %r55, 9;
>>          @%r72   bra     $L3;
>>   $L2:
>> -               bar.sync        1;
>>          // joining 2;
>>   $L5:
>> +               bar.sync        1;
>>          // join 2;
>>          ret;
>>   }
>>
>>
> 
> At -O0, yes.
> 
> At -O2, we have:
> ...
>  diff -u -a 1 2
> --- 1   2018-03-19 14:13:44.074834552 +0100
> +++ 2   2018-03-19 14:15:06.075301168 +0100
> @@ -42,20 +42,20 @@
>  st.u64 [%r32],%r25;
>  $L7:
>  $L6:
> -@ %r33 bra $L5;
>  // forked 2;
>  bar.sync 0;
> +@ %r33 bra $L5;
>  cvta.shared.u64 %r31,__worker_bcast;
>  ld.u64 %r25,[%r31];
>  mov.u32 %r24,%tid.y;
>  setp.le.s32 %r26,%r24,9;
>  @ %r26 bra $L2;
>  $L3:
> -bar.sync 1;
>  // joining 2;
>  $L5:
> -@ %r34 bra.uni $L8;
> +bar.sync 1;
>  @ %r33 bra $L9;
> +@ %r34 bra.uni $L8;
>  // join 2;
>  $L9:
>  $L8:
> ...
> 
> Note that this changes ordering of the vector-neutering jump and
> worker-neutering jump at the end. In principle, this should not be
> harmful, but it violates the invariant that vector-neutering
> branch-around code should be as short-lived as possible. So, this needs
> to be fixed.
> 
> I've found this issue by adding verification of the neutering, as
> attached below.

ACK, thanks. I'll take a closer look at this.

Is your patch purely for debugging, or are you planning on committing it
to og7 and trunk?

Cesar