<div dir="ltr">Hi David,<div><br></div><div>Just to make sure we are on the same page. We are talking about nvcc's behavior here. LLVM does the right thing to me (i.e. not duplicating) on this example. I have no idea how nvcc is implemented. </div></div><div class="gmail_extra"><br><div class="gmail_quote">On Fri, Aug 21, 2015 at 2:00 PM, David Majnemer <span dir="ltr"><<a href="mailto:david.majnemer@gmail.com" target="_blank">david.majnemer@gmail.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr">Perhaps it is semantics preserving so long as the __syncthreads callsite is marked noduplicate?<div><br></div><div><a href="https://github.com/llvm-mirror/llvm/blob/896f064a4900458e3fb245ad3f6fc9e7a3d8c8cd/lib/Analysis/InlineCost.cpp#L1284" target="_blank">https://github.com/llvm-mirror/llvm/blob/896f064a4900458e3fb245ad3f6fc9e7a3d8c8cd/lib/Analysis/InlineCost.cpp#L1284</a><br></div><div class="gmail_extra"><br><div class="gmail_quote"><div><div class="h5">On Fri, Aug 21, 2015 at 1:56 PM, Justin Holewinski via llvm-dev <span dir="ltr"><<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>></span> wrote:<br></div></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div><div class="h5">That’s an interesting case.  AFAIK, inlining should be restricted here for the reason you mention.  Inlining should only be valid if it doesn’t duplicate the barrier.  Which nvcc shows this behavior?<br>
<br>
Adding Vinod and Yuan for comment.<br>
<div><div><br>
> On Aug 21, 2015, at 4:24 PM, Jingyue Wu <<a href="mailto:jingyue@google.com" target="_blank">jingyue@google.com</a>> wrote:<br>
><br>
> Hi Justin,<br>
><br>
> Is a compiler allowed to inline a function that calls __syncthreads? I saw nvcc does that, but not sure it's valid though. For example,<br>
><br>
> void foo() {<br>
>   __syncthreads();<br>
> }<br>
><br>
> if (threadIdx.x % 2 == 0) {<br>
>   ...<br>
>   foo();<br>
> } else {<br>
>   ...<br>
>   foo();<br>
> }<br>
><br>
> Before inlining, all threads meet at one __syncthreads(). After inlining<br>
><br>
> if (threadIdx.x % 2 == 0) {<br>
>   ...<br>
>   __syncthreads();<br>
> } else {<br>
>   ...<br>
>   __syncthreads();<br>
> }<br>
><br>
> The __syncthreads call is duplicated, and it's no longer guaranteed that all threads can meet one __syncthreads().<br>
><br>
> Any thoughts?<br>
><br>
> Jingyue<br>
<br>
<br>
</div></div>-----------------------------------------------------------------------------------<br>
This email message is for the sole use of the intended recipient(s) and may contain<br>
confidential information.  Any unauthorized review, use, disclosure or distribution<br>
is prohibited.  If you are not the intended recipient, please contact the sender by<br>
reply email and destroy all copies of the original message.<br>
-----------------------------------------------------------------------------------<br></div></div>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</blockquote></div><br></div></div>
</blockquote></div><br></div>