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....@amd.com]
Sent: 20 October 2016 17:59
To: Ettore Speziale; reviews+d25343+public+a10e9553b0fc8...@reviews.llvm.org
Cc: alexey.ba...@intel.com; Anastasia Stulova; aaron.ball...@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.ett...@gmail.com]
Sent: Thursday, October 20, 2016 11:42 AM
To: reviews+d25343+public+a10e9553b0fc8...@reviews.llvm.org; Liu, Yaxun (Sam)
<yaxun....@amd.com>
Cc: Ettore Speziale <speziale.ett...@gmail.com>; alexey.ba...@intel.com;
anastasia.stul...@arm.com; aaron.ball...@gmail.com; Clang Commits
<cfe-commits@lists.llvm.org>; Sumner, Brian <brian.sum...@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.ett...@gmail.com
espezi...@apple.com
--------------------------------------------------
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
