<br><br>On Friday, August 22, 2014, Tom Stellard <<a href="mailto:tom@stellard.net">tom@stellard.net</a>> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
On Fri, Aug 22, 2014 at 02:33:43AM +0200, Hilloulin Damien wrote:<br>
>  This patch is a first implementation of the newly added intrinsics<br>
>  for barriers/memory fences on Evergreen. For ultra-simplicity, every<br>
>  intrinsic is lowered to a barrier with no fence. But for real fences, we<br>
>  would need to place WAIT_ACK instructions and transform surrounding<br>
>  read/write instructions to make ACK.<br>
><br>
<br>
As I mentioned in the libclc review, the first step should be<br>
to implement generic versions of memfence which just calls barrier.<br>
That way we won't need to make any changes in the backend.</blockquote><div>It is not correct to use barrier at all to implement mem_fence. barrier requires all work items to reach the barrier, and mem_fence does not. s_barrier also does not imply completed memory operations</div>
<div> </div><div>If (get_local_id(0) == 0)</div><div>{</div><div>    gptr = 1;</div><div>    mem_fence(CLK_GLOBAL_MEM_FENCE);</div><div>}</div><div><br></div><div>Is a valid use of a fence but not a barrier<span></span></div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<br>
Once the generic implementations have been committed, we can go through<br>
and optimize each function one at a time.  This will make for a cleaner<br>
set of changes.<br>
<br>
-Tom<br>
<br>
> Signed-off-by: Damien Hilloulin <<a href="javascript:;" onclick="_e(event, 'cvml', 'damien.hilloulin@supelec.fr')">damien.hilloulin@supelec.fr</a>><br>
> ---<br>
>  lib/Target/R600/EvergreenInstructions.td | 69<br>
> +++++++++++++++++++++++++++++++-<br>
>  1 file changed, 68 insertions(+), 1 deletion(-)<br>
><br>
> diff --git a/lib/Target/R600/EvergreenInstructions.td<br>
> b/lib/Target/R600/EvergreenInstructions.td<br>
> index a83567a..c8f90ce 100644<br>
> --- a/lib/Target/R600/EvergreenInstructions.td<br>
> +++ b/lib/Target/R600/EvergreenInstructions.td<br>
> @@ -358,8 +358,12 @@ def FLT_TO_UINT_eg : FLT_TO_UINT_Common<0x9A> {<br>
><br>
>  def UINT_TO_FLT_eg : UINT_TO_FLT_Common<0x9C>;<br>
><br>
> +//===----------------------------------------------------------------------===//<br>
> +// SYnchronization instructions<br>
> +//===----------------------------------------------------------------------===//<br>
> +<br>
>  def GROUP_BARRIER : InstR600 <<br>
> -    (outs), (ins), "  GROUP_BARRIER", [(int_AMDGPU_barrier_local),<br>
> (int_AMDGPU_barrier_global)], AnyALU>,<br>
> +    (outs), (ins), "  GROUP_BARRIER", [], AnyALU>,<br>
>      R600ALU_Word0,<br>
>      R600ALU_Word1_OP2 <0x54> {<br>
><br>
> @@ -389,10 +393,73 @@ def GROUP_BARRIER : InstR600 <<br>
>  }<br>
><br>
>  def : Pat <<br>
> +    (int_AMDGPU_barrier_nofence),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +// XXX: the following patterns in the section are stubs.<br>
> +// We should take care of inserting WAIT_ACK and modifying the<br>
> +// read/writes instructions before the barrier and in the loop.<br>
> +def : Pat <<br>
> +    (int_AMDGPU_barrier_local),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
>      (int_AMDGPU_barrier_global),<br>
>      (GROUP_BARRIER)<br>
>  >;<br>
><br>
> +def : Pat <<br>
> +    (int_AMDGPU_barrier_localglobal),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_mem_fence_local),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_mem_fence_global),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_mem_fence_localglobal),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_read_mem_fence_local),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_read_mem_fence_global),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_read_mem_fence_localglobal),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_write_mem_fence_local),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_write_mem_fence_global),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
> +<br>
> +def : Pat <<br>
> +    (int_AMDGPU_write_mem_fence_localglobal),<br>
> +    (GROUP_BARRIER)<br>
> +>;<br>
>  //===----------------------------------------------------------------------===//<br>
>  // LDS Instructions<br>
>  //===----------------------------------------------------------------------===//<br>
> --<br>
> 1.9.1<br>
><br>
> _______________________________________________<br>
> llvm-commits mailing list<br>
> <a href="javascript:;" onclick="_e(event, 'cvml', 'llvm-commits@cs.uiuc.edu')">llvm-commits@cs.uiuc.edu</a><br>
> <a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits</a><br>
_______________________________________________<br>
llvm-commits mailing list<br>
<a href="javascript:;" onclick="_e(event, 'cvml', 'llvm-commits@cs.uiuc.edu')">llvm-commits@cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits</a><br>
</blockquote>