[PATCH] D25343: [OpenCL] Mark group functions as convergent in opencl-c.h

[Liu, Yaxun (Sam) via cfe-commits]

+ Tom Matt

Thanks Ettore.
 
I think OpenCL is subject to the same issue, and noduplicate does not help either.

Basically if a function A directly or indirectly calls a convergent function e.g. barrier, function A itself must also be marked as convergent, otherwise optimization passes may transform a convergent call of A into multiple divergent calls of A.

That means if we only know the declaration of a function, we have to assume it is convergent since in its body it may call a convergent function.

I think probably OpenCL should take the same approach, i.e., mark all functions as convergent, then let Transforms/IPO/FunctionAttrs.cpp to remove unnecessary convergent attribute.

Sam

-----Original Message-----
From: Ettore Speziale [mailto:speziale.ettore at gmail.com] 
Sent: Thursday, October 20, 2016 11:42 AM
To: reviews+D25343+public+a10e9553b0fc895f at reviews.llvm.org; Liu, Yaxun (Sam) <Yaxun.Liu at amd.com>
Cc: Ettore Speziale <speziale.ettore at gmail.com>; alexey.bader at intel.com; anastasia.stulova at arm.com; aaron.ballman at gmail.com; Clang Commits <cfe-commits at lists.llvm.org>; Sumner, Brian <Brian.Sumner at amd.com>
Subject: Re: [PATCH] D25343: [OpenCL] Mark group functions as convergent in opencl-c.h

Hello guys,

>> Should we deprecate noduplicate then as convergent should cover both use cases for OpenCL I believe? As far as I understand noduplicate was added specifically for SPMD use cases...
> 
> noduplicate has different semantics than convergent. Although it is proposed for SPMD originally, it could be used by non-SPMD programs to forbid duplicate of functions. There may be applications using this attribute.
> 
> I would suggest to leave this question for future. Probably ask llvm-dev first since the attribute is also in LLVM.

I just want to clarify why I withdraw the convergent patch I initially submitted some time ago. It has a problem when dealing with multiple modules. Consider the following example:

int foo(void) __attribute__((pure));

int bar(int x) {
  int y = foo();
  if (x)
    return y;
  return 0;
}   

I’ve just marked foo with the pure attribute to mark the function readonly in LLVM IR. Given that IR, the IR sinking pass pushes the foo call site into the then branch:

; Function Attrs: nounwind readonly ssp uwtable define i32 @bar(i32) #0 {
  %2 = icmp eq i32 %0, 0
  br i1 %2, label %5, label %3

; <label>:3                                       ; preds = %1
  %4 = tail call i32 @foo() #2
  br label %5

; <label>:5                                       ; preds = %1, %3
  %6 = phi i32 [ %4, %3 ], [ 0, %1 ]
  ret i32 %6
}

; Function Attrs: nounwind readonly
declare i32 @foo() #1

This is kind of dangerous, as we do not know what is inside foo — i.e. it might contains a convergent call.

If I understand correctly, the CUDA guys solved the problem in two steps. At CodeGen time all the device function calls are marked convergent:

  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
    // Conservatively, mark all functions and calls in CUDA as convergent
    // (meaning, they may call an intrinsically convergent op, such as
    // __syncthreads(), and so can't have certain optimizations applied around
    // them).  LLVM will remove this attribute where it safely can.
    FuncAttrs.addAttribute(llvm::Attribute::Convergent);

Then LLVM function attribute pass — lib/Transforms/IPO/FunctionAttrs.cpp — remove the unnecessary convergent attributes starting from the leaf nodes  — i.e. external calls.

Provide that intrinsics are correctly marked convergent only when needed, that allow to get rid of the unnecessary convergent attributes.

Since you are introducing an explicit convergent attribute it seems that OpenCL is requiring the developers to explicitly mark the functions that might contain convergent function calls with the convergent attribute. Am I understand correctly?

Thanks

--------------------------------------------------
Ettore Speziale — Compiler Engineer
speziale.ettore at gmail.com
espeziale at apple.com
--------------------------------------------------