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

[Anastasia Stulova via cfe-commits]

Hi Ettore,

As far as I understand the whole problem is that the optimized functions are marked by __attribute__((pure)). If the attribute is removed from your example, we get LLVM dump preserving correctness:

define i32 @bar(i32 %x) local_unnamed_addr #0 {
entry:
  %call = tail call i32 @foo() #2
  %tobool = icmp eq i32 %x, 0
  %.call = select i1 %tobool, i32 0, i32 %call
  ret i32 %.call
}

Because by default a function without a definition is considered to have side effects. If a function has definition at some point and either it's mapped to llvm intrintic with convergent attribute or it calls an intrinsic with a convergent attribute in LLVM then LLVM should make sure not to perform any illegal optimisation. I am not an expert in LLVM however. @Tom or anyone else experienced, would you be able to comment whether this is the case?

However, if I understand it well that the motivation for adding the convergent attribute to Clang could be to allow some functions (which now have to be considered side effect functions) to be better optimised by marking them as convergent instead of leaving as side effect? Convergent functions can be optimised within its BB as soon as CFG equivalence is preserved for the whole caller function.

Appending convergent attribute to every function in Clang CodeGen seems overly conservative and can prevent many optimisation. I am not sure how good is that LLVM pass removing the attribute. Surely removing the attribute will be difficult for the functions with definitions in separate compilation units. Do we have a good case for adding the convergent attribute to the frontend now?

I am not sure we have a use case for the noduplicate as well. To me it seems that CFG equivalence sematic of convergent attribute would cover noduplicate use cases too.

Cheers,
Anastasia

-----Original Message-----
From: Liu, Yaxun (Sam) [mailto:Yaxun.Liu at amd.com] 
Sent: 20 October 2016 17:59
To: Ettore Speziale; reviews+D25343+public+a10e9553b0fc895f at reviews.llvm.org
Cc: alexey.bader at intel.com; Anastasia Stulova; aaron.ballman at gmail.com; Clang Commits; Sumner, Brian; Stellard, Thomas; Arsenault, Matthew
Subject: RE: [PATCH] D25343: [OpenCL] Mark group functions as convergent in opencl-c.h

+ 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
--------------------------------------------------