<div dir="ltr"><br><div class="gmail_extra"><br><div class="gmail_quote">On Thu, Jul 21, 2016 at 12:18 PM, Jan Vesely <span dir="ltr"><<a href="mailto:jan.vesely@rutgers.edu" target="_blank">jan.vesely@rutgers.edu</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div class="HOEnZb"><div class="h5">On Fri, 2016-07-15 at 11:59 -0400, Jan Vesely wrote:<br>
> Signed-off-by: Jan Vesely <<a href="mailto:jan.vesely@rutgers.edu">jan.vesely@rutgers.edu</a>><br>
> ---<br>
> I'm pretty sure barrier implementation is busted, as it's noop<br>
> without CLK_LOCAL_MEM_FENCE.<br>
><br>
>  ptx-nvidiacl/lib/synchronization/<a href="http://barrier.cl" rel="noreferrer" target="_blank">barrier.cl</a> | 2 +-<br>
>  ptx-nvidiacl/lib/workitem/<a href="http://get_group_id.cl" rel="noreferrer" target="_blank">get_group_id.cl</a>   | 6 +++---<br>
>  ptx-nvidiacl/lib/workitem/<a href="http://get_local_id.cl" rel="noreferrer" target="_blank">get_local_id.cl</a>   | 6 +++---<br>
>  ptx-nvidiacl/lib/workitem/<a href="http://get_local_size.cl" rel="noreferrer" target="_blank">get_local_size.cl</a> | 6 +++---<br>
>  ptx-nvidiacl/lib/workitem/<a href="http://get_num_groups.cl" rel="noreferrer" target="_blank">get_num_groups.cl</a> | 6 +++---<br>
>  5 files changed, 13 insertions(+), 13 deletions(-)<br>
><br>
> diff --git a/ptx-nvidiacl/lib/synchronization/<a href="http://barrier.cl" rel="noreferrer" target="_blank">barrier.cl</a> b/ptx-<br>
> nvidiacl/lib/synchronization/<a href="http://barrier.cl" rel="noreferrer" target="_blank">barrier.cl</a><br>
> index fb36c26..88e1493 100644<br>
> --- a/ptx-nvidiacl/lib/synchronization/<a href="http://barrier.cl" rel="noreferrer" target="_blank">barrier.cl</a><br>
> +++ b/ptx-nvidiacl/lib/synchronization/<a href="http://barrier.cl" rel="noreferrer" target="_blank">barrier.cl</a><br>
> @@ -2,7 +2,7 @@<br>
>  <br>
>  _CLC_DEF void barrier(cl_mem_fence_flags flags) {<br>
>    if (flags & CLK_LOCAL_MEM_FENCE) {<br>
> -    __builtin_ptx_bar_sync(0);<br>
> +    __syncthreads();<br>
>    }<br>
>  }<br>
>  <br>
> diff --git a/ptx-nvidiacl/lib/workitem/<a href="http://get_group_id.cl" rel="noreferrer" target="_blank">get_group_id.cl</a> b/ptx-<br>
> nvidiacl/lib/workitem/<a href="http://get_group_id.cl" rel="noreferrer" target="_blank">get_group_id.cl</a><br>
> index 2b35b4e..dbc4784 100644<br>
> --- a/ptx-nvidiacl/lib/workitem/<a href="http://get_group_id.cl" rel="noreferrer" target="_blank">get_group_id.cl</a><br>
> +++ b/ptx-nvidiacl/lib/workitem/<a href="http://get_group_id.cl" rel="noreferrer" target="_blank">get_group_id.cl</a><br>
> @@ -2,9 +2,9 @@<br>
>  <br>
>  _CLC_DEF size_t get_group_id(uint dim) {<br>
>    switch (dim) {<br>
> -  case 0:  return __builtin_ptx_read_ctaid_x();<br>
> -  case 1:  return __builtin_ptx_read_ctaid_y();<br>
> -  case 2:  return __builtin_ptx_read_ctaid_z();<br>
> +  case 0:  return __nvvm_read_ptx_sreg_ctaid_x();<br>
> +  case 1:  return __nvvm_read_ptx_sreg_ctaid_y();<br>
> +  case 2:  return __nvvm_read_ptx_sreg_ctaid_z();<br>
>    default: return 0;<br>
>    }<br>
>  }<br>
> diff --git a/ptx-nvidiacl/lib/workitem/<a href="http://get_local_id.cl" rel="noreferrer" target="_blank">get_local_id.cl</a> b/ptx-<br>
> nvidiacl/lib/workitem/<a href="http://get_local_id.cl" rel="noreferrer" target="_blank">get_local_id.cl</a><br>
> index f0cfdc0..f31581a 100644<br>
> --- a/ptx-nvidiacl/lib/workitem/<a href="http://get_local_id.cl" rel="noreferrer" target="_blank">get_local_id.cl</a><br>
> +++ b/ptx-nvidiacl/lib/workitem/<a href="http://get_local_id.cl" rel="noreferrer" target="_blank">get_local_id.cl</a><br>
> @@ -2,9 +2,9 @@<br>
>  <br>
>  _CLC_DEF size_t get_local_id(uint dim) {<br>
>    switch (dim) {<br>
> -  case 0:  return __builtin_ptx_read_tid_x();<br>
> -  case 1:  return __builtin_ptx_read_tid_y();<br>
> -  case 2:  return __builtin_ptx_read_tid_z();<br>
> +  case 0:  return __nvvm_read_ptx_sreg_tid_x();<br>
> +  case 1:  return __nvvm_read_ptx_sreg_tid_y();<br>
> +  case 2:  return __nvvm_read_ptx_sreg_tid_z();<br>
>    default: return 0;<br>
>    }<br>
>  }<br>
> diff --git a/ptx-nvidiacl/lib/workitem/<a href="http://get_local_size.cl" rel="noreferrer" target="_blank">get_local_size.cl</a> b/ptx-<br>
> nvidiacl/lib/workitem/<a href="http://get_local_size.cl" rel="noreferrer" target="_blank">get_local_size.cl</a><br>
> index c3f5425..d00b0d6 100644<br>
> --- a/ptx-nvidiacl/lib/workitem/<a href="http://get_local_size.cl" rel="noreferrer" target="_blank">get_local_size.cl</a><br>
> +++ b/ptx-nvidiacl/lib/workitem/<a href="http://get_local_size.cl" rel="noreferrer" target="_blank">get_local_size.cl</a><br>
> @@ -2,9 +2,9 @@<br>
>  <br>
>  _CLC_DEF size_t get_local_size(uint dim) {<br>
>    switch (dim) {<br>
> -  case 0:  return __builtin_ptx_read_ntid_x();<br>
> -  case 1:  return __builtin_ptx_read_ntid_y();<br>
> -  case 2:  return __builtin_ptx_read_ntid_z();<br>
> +  case 0:  return __nvvm_read_ptx_sreg_ntid_x();<br>
> +  case 1:  return __nvvm_read_ptx_sreg_ntid_y();<br>
> +  case 2:  return __nvvm_read_ptx_sreg_ntid_z();<br>
>    default: return 0;<br>
>    }<br>
>  }<br>
> diff --git a/ptx-nvidiacl/lib/workitem/<a href="http://get_num_groups.cl" rel="noreferrer" target="_blank">get_num_groups.cl</a> b/ptx-<br>
> nvidiacl/lib/workitem/<a href="http://get_num_groups.cl" rel="noreferrer" target="_blank">get_num_groups.cl</a><br>
> index 90bdc2e..d7abf3f 100644<br>
> --- a/ptx-nvidiacl/lib/workitem/<a href="http://get_num_groups.cl" rel="noreferrer" target="_blank">get_num_groups.cl</a><br>
> +++ b/ptx-nvidiacl/lib/workitem/<a href="http://get_num_groups.cl" rel="noreferrer" target="_blank">get_num_groups.cl</a><br>
> @@ -2,9 +2,9 @@<br>
>  <br>
>  _CLC_DEF size_t get_num_groups(uint dim) {<br>
>    switch (dim) {<br>
> -  case 0:  return __builtin_ptx_read_nctaid_x();<br>
> -  case 1:  return __builtin_ptx_read_nctaid_y();<br>
> -  case 2:  return __builtin_ptx_read_nctaid_z();<br>
> +  case 0:  return __nvvm_read_ptx_sreg_nctaid_x();<br>
> +  case 1:  return __nvvm_read_ptx_sreg_nctaid_y();<br>
> +  case 2:  return __nvvm_read_ptx_sreg_nctaid_z();<br>
>    default: return 0;<br>
>    }<br>
>  }<br>
<br>
</div></div>ping, this is needed to fix nvptx build.<span class="HOEnZb"><font color="#888888"><br></font></span></blockquote><div><br></div><div>I don't know PTX well enough to provide a proper review, but it looks sane enough to me, and I can confirm that it fixes the build for me (and that it was broken before).<br><br></div><div>How about an Acked-By: Aaron Watry <<a href="mailto:awatry@gmail.com">awatry@gmail.com</a>><br></div><div><br></div><div>--Aaron<br></div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><span class="HOEnZb"><font color="#888888">
--<br>
Jan Vesely <<a href="mailto:jan.vesely@rutgers.edu">jan.vesely@rutgers.edu</a>></font></span></blockquote></div><br></div></div>