<div dir="ltr">Hello all (sorry for the long post),<br><br>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: <div>
<br></div><div>

<b>  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. </b></div>
<b>  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.<br>  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. </b><div>
<br></div><div>So if we have the following SPIR:<br><br>entry:<br>  %tid = tail call i32 @get_global_id(i32 0)<br>  %cmp = icmp sgt i32 %tid, 25<br>  br i1 %cmp, label %if, label %exit<br>
<br>if:<br>  ...<br>  tail call void @barrier(i32 3)<br>  ...<br>  br label %exit<br>
<br>exit:<br>  ret void<br><br></div><div>The barrier call is thread dependent and violate uniformness. However if we have another SPIR:</div><div><br></div><div>entry:<br>  %tid = tail call i32 @get_global_id(i32 0)<br>
  %cmp = icmp sgt i32 %tid, 25<br>  br i1 %cmp, label %if, label %else<br><br>if:<br>  ...<br>  tail call void @barrier(i32 3)<br>  ...<br>  br label %exit</div><div><br></div><div>else:</div><div>  ...</div><div>  tail call void @barrier(i32 3)</div>
<div>  ...</div><div>  br label %exit<br><br>exit:<br>  ret void<br></div><div><br></div><div>The barrier function would be executed regardless of the value returned by get_global_id so it is thread independent.</div><div>
<br></div><div>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.</div>
<div><br></div><div>Any suggestion and keyword is welcomed and appreciated :)</div><div><br></div></div><div class="gmail_extra"><br><br><div class="gmail_quote">2014-05-20 23:26 GMT+08:00 Renato Golin <span dir="ltr"><<a href="mailto:renato.golin@linaro.org" target="_blank">renato.golin@linaro.org</a>></span>:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div class="">On 20 May 2014 16:08, RICHARD STUCKEY <<a href="mailto:richard.stuckey@virgin.net">richard.stuckey@virgin.net</a>> wrote:<br>

> Consider a function which contains an infinite loop: any algorithm which<br>
> could determine whether that function is called or not would effectively be<br>
> an algorithm that could determine whether the program containing that<br>
> function halts or not.  Equally, deciding whether the function contains an<br>
> infinite loop or not is itself an instance of the Halting Problem.<br>
<br>
</div>A more conservative approach, bailing out on any loops or<br>
complication, would scan the BB flow graph in between any two nodes<br>
(between fully dominating / fully dominated) would return { no, maybe,<br>
yes }. I'd guess that it'd return maybe much more often that yes/no<br>
for the interesting cases, limiting the possibilities of this<br>
approach.<br>
<br>
Do you have a concrete use in mind? DCE for the "no" cases? Hoisting<br>
for the "yes" cases? Wouldn't you have to track the arguments as well?<br>
That can quickly get out of hand...<br>
<br>
cheers,<br>
--renato<br>
</blockquote></div><br></div>