Renato Golin
2014-May-20 15:26 UTC
[LLVMdev] How to decide whether a function is executed or not
On 20 May 2014 16:08, RICHARD STUCKEY <richard.stuckey at virgin.net> wrote:> Consider a function which contains an infinite loop: any algorithm which > could determine whether that function is called or not would effectively be > an algorithm that could determine whether the program containing that > function halts or not. Equally, deciding whether the function contains an > infinite loop or not is itself an instance of the Halting Problem.A more conservative approach, bailing out on any loops or complication, would scan the BB flow graph in between any two nodes (between fully dominating / fully dominated) would return { no, maybe, yes }. I'd guess that it'd return maybe much more often that yes/no for the interesting cases, limiting the possibilities of this approach. Do you have a concrete use in mind? DCE for the "no" cases? Hoisting for the "yes" cases? Wouldn't you have to track the arguments as well? That can quickly get out of hand... cheers, --renato
lyh.kernel
2014-May-22 19:01 UTC
[LLVMdev] How to decide whether a function is executed or not
Hello all (sorry for the long post), Sorry for the misleading of the previous mail. In fact what I want to implement is a uniformness checking pass of OpenCL barrier. Here is the description of barrier function: * All work-items in a work-group executing the kernel on a processor must execute this function before any are allowed to continue execution beyond the barrier. This function must be encountered by all work-items in a work-group executing the kernel. * * If barrier is inside a conditional statement, then all work-items must enter the conditional if any work-item enters the conditional statement and executes the barrier. If barrier is inside a loop, all work-items must execute the barrier for each iteration of the loop before any are allowed to continue execution beyond the barrier. * So if we have the following SPIR: entry: %tid = tail call i32 @get_global_id(i32 0) %cmp = icmp sgt i32 %tid, 25 br i1 %cmp, label %if, label %exit if: ... tail call void @barrier(i32 3) ... br label %exit exit: ret void The barrier call is thread dependent and violate uniformness. However if we have another SPIR: entry: %tid = tail call i32 @get_global_id(i32 0) %cmp = icmp sgt i32 %tid, 25 br i1 %cmp, label %if, label %else if: ... tail call void @barrier(i32 3) ... br label %exit else: ... tail call void @barrier(i32 3) ... br label %exit exit: ret void The barrier function would be executed regardless of the value returned by get_global_id so it is thread independent. I have implement a LLVM pass to detect control dependency between "tail call void @barrier(i32 3)" and "%tid = tail call i32 @get_global_id(i32 0)". But the pass is unuseful for the second example. Any suggestion and keyword is welcomed and appreciated :) 2014-05-20 23:26 GMT+08:00 Renato Golin <renato.golin at linaro.org>:> On 20 May 2014 16:08, RICHARD STUCKEY <richard.stuckey at virgin.net> wrote: > > Consider a function which contains an infinite loop: any algorithm which > > could determine whether that function is called or not would effectively > be > > an algorithm that could determine whether the program containing that > > function halts or not. Equally, deciding whether the function contains > an > > infinite loop or not is itself an instance of the Halting Problem. > > A more conservative approach, bailing out on any loops or > complication, would scan the BB flow graph in between any two nodes > (between fully dominating / fully dominated) would return { no, maybe, > yes }. I'd guess that it'd return maybe much more often that yes/no > for the interesting cases, limiting the possibilities of this > approach. > > Do you have a concrete use in mind? DCE for the "no" cases? Hoisting > for the "yes" cases? Wouldn't you have to track the arguments as well? > That can quickly get out of hand... > > cheers, > --renato >-------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20140523/6c0e3317/attachment.html>
Peter Collingbourne
2014-May-22 21:06 UTC
[LLVMdev] How to decide whether a function is executed or not
You might want to look at this paper: http://www.doc.ic.ac.uk/~afd/homepages/papers/pdfs/2013/ESOP.pdf It defines a notion of barrier divergence and lays out a scheme for showing that a program is free from it. There is also an implementation that uses a theorem prover to show barrier divergence freedom (among other things). Peter On Fri, May 23, 2014 at 03:01:21AM +0800, lyh.kernel wrote:> Hello all (sorry for the long post), > > Sorry for the misleading of the previous mail. In fact what I want to > implement is a uniformness checking pass of OpenCL barrier. Here is the > description of barrier function: > > * All work-items in a work-group executing the kernel on a processor must > execute this function before any are allowed to continue execution beyond > the barrier. This function must be encountered by all work-items in a > work-group executing the kernel. * > > * If barrier is inside a conditional statement, then all work-items must > enter the conditional if any work-item enters the conditional statement and > executes the barrier. If barrier is inside a loop, all work-items must > execute the barrier for each iteration of the loop before any are allowed > to continue execution beyond the barrier. * > > So if we have the following SPIR: > > entry: > %tid = tail call i32 @get_global_id(i32 0) > %cmp = icmp sgt i32 %tid, 25 > br i1 %cmp, label %if, label %exit > > if: > ... > tail call void @barrier(i32 3) > ... > br label %exit > > exit: > ret void > > The barrier call is thread dependent and violate uniformness. However if we > have another SPIR: > > entry: > %tid = tail call i32 @get_global_id(i32 0) > %cmp = icmp sgt i32 %tid, 25 > br i1 %cmp, label %if, label %else > > if: > ... > tail call void @barrier(i32 3) > ... > br label %exit > > else: > ... > tail call void @barrier(i32 3) > ... > br label %exit > > exit: > ret void > > The barrier function would be executed regardless of the value returned by > get_global_id so it is thread independent. > > I have implement a LLVM pass to detect control dependency between "tail > call void @barrier(i32 3)" and "%tid = tail call i32 @get_global_id(i32 > 0)". But the pass is unuseful for the second example. > > Any suggestion and keyword is welcomed and appreciated :) > > > > 2014-05-20 23:26 GMT+08:00 Renato Golin <renato.golin at linaro.org>: > > > On 20 May 2014 16:08, RICHARD STUCKEY <richard.stuckey at virgin.net> wrote: > > > Consider a function which contains an infinite loop: any algorithm which > > > could determine whether that function is called or not would effectively > > be > > > an algorithm that could determine whether the program containing that > > > function halts or not. Equally, deciding whether the function contains > > an > > > infinite loop or not is itself an instance of the Halting Problem. > > > > A more conservative approach, bailing out on any loops or > > complication, would scan the BB flow graph in between any two nodes > > (between fully dominating / fully dominated) would return { no, maybe, > > yes }. I'd guess that it'd return maybe much more often that yes/no > > for the interesting cases, limiting the possibilities of this > > approach. > > > > Do you have a concrete use in mind? DCE for the "no" cases? Hoisting > > for the "yes" cases? Wouldn't you have to track the arguments as well? > > That can quickly get out of hand... > > > > cheers, > > --renato > >> _______________________________________________ > LLVM Developers mailing list > LLVMdev at cs.uiuc.edu http://llvm.cs.uiuc.edu > http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev-- Peter