[llvm-dev] [CUDA/NVPTX] is inlining __syncthreads allowed?

[Jingyue Wu via llvm-dev]

Hi David,

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.

On Fri, Aug 21, 2015 at 2:00 PM, David Majnemer <david.majnemer at gmail.com>
wrote:

> Perhaps it is semantics preserving so long as the __syncthreads callsite
> is marked noduplicate?
>
>
> https://github.com/llvm-mirror/llvm/blob/896f064a4900458e3fb245ad3f6fc9e7a3d8c8cd/lib/Analysis/InlineCost.cpp#L1284
>
> On Fri, Aug 21, 2015 at 1:56 PM, Justin Holewinski via llvm-dev <
> llvm-dev at lists.llvm.org> wrote:
>
>> 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?
>>
>> Adding Vinod and Yuan for comment.
>>
>> > On Aug 21, 2015, at 4:24 PM, Jingyue Wu <jingyue at google.com> wrote:
>> >
>> > Hi Justin,
>> >
>> > 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,
>> >
>> > void foo() {
>> >   __syncthreads();
>> > }
>> >
>> > if (threadIdx.x % 2 == 0) {
>> >   ...
>> >   foo();
>> > } else {
>> >   ...
>> >   foo();
>> > }
>> >
>> > Before inlining, all threads meet at one __syncthreads(). After inlining
>> >
>> > if (threadIdx.x % 2 == 0) {
>> >   ...
>> >   __syncthreads();
>> > } else {
>> >   ...
>> >   __syncthreads();
>> > }
>> >
>> > The __syncthreads call is duplicated, and it's no longer guaranteed
>> that all threads can meet one __syncthreads().
>> >
>> > Any thoughts?
>> >
>> > Jingyue
>>
>>
>>
>> -----------------------------------------------------------------------------------
>> This email message is for the sole use of the intended recipient(s) and
>> may contain
>> confidential information.  Any unauthorized review, use, disclosure or
>> distribution
>> is prohibited.  If you are not the intended recipient, please contact the
>> sender by
>> reply email and destroy all copies of the original message.
>>
>> -----------------------------------------------------------------------------------
>> _______________________________________________
>> LLVM Developers mailing list
>> llvm-dev at lists.llvm.org
>> http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
>>
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20150821/992cd546/attachment.html>