[LLVMdev] How to decide whether a function is executed or not

[Peter Collingbourne]

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