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

[lyh.kernel]

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>