[Libclc-dev] [PATCH 1/1] ptx: Fix builtin names after clang r274770

[Jan Vesely via Libclc-dev]

On Fri, 2016-07-15 at 11:59 -0400, Jan Vesely wrote:
> Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
> ---
> I'm pretty sure barrier implementation is busted, as it's noop
> without CLK_LOCAL_MEM_FENCE.
> 
>  ptx-nvidiacl/lib/synchronization/barrier.cl | 2 +-
>  ptx-nvidiacl/lib/workitem/get_group_id.cl   | 6 +++---
>  ptx-nvidiacl/lib/workitem/get_local_id.cl   | 6 +++---
>  ptx-nvidiacl/lib/workitem/get_local_size.cl | 6 +++---
>  ptx-nvidiacl/lib/workitem/get_num_groups.cl | 6 +++---
>  5 files changed, 13 insertions(+), 13 deletions(-)
> 
> diff --git a/ptx-nvidiacl/lib/synchronization/barrier.cl b/ptx-
> nvidiacl/lib/synchronization/barrier.cl
> index fb36c26..88e1493 100644
> --- a/ptx-nvidiacl/lib/synchronization/barrier.cl
> +++ b/ptx-nvidiacl/lib/synchronization/barrier.cl
> @@ -2,7 +2,7 @@
>  
>  _CLC_DEF void barrier(cl_mem_fence_flags flags) {
>    if (flags & CLK_LOCAL_MEM_FENCE) {
> -    __builtin_ptx_bar_sync(0);
> +    __syncthreads();
>    }
>  }
>  
> diff --git a/ptx-nvidiacl/lib/workitem/get_group_id.cl b/ptx-
> nvidiacl/lib/workitem/get_group_id.cl
> index 2b35b4e..dbc4784 100644
> --- a/ptx-nvidiacl/lib/workitem/get_group_id.cl
> +++ b/ptx-nvidiacl/lib/workitem/get_group_id.cl
> @@ -2,9 +2,9 @@
>  
>  _CLC_DEF size_t get_group_id(uint dim) {
>    switch (dim) {
> -  case 0:  return __builtin_ptx_read_ctaid_x();
> -  case 1:  return __builtin_ptx_read_ctaid_y();
> -  case 2:  return __builtin_ptx_read_ctaid_z();
> +  case 0:  return __nvvm_read_ptx_sreg_ctaid_x();
> +  case 1:  return __nvvm_read_ptx_sreg_ctaid_y();
> +  case 2:  return __nvvm_read_ptx_sreg_ctaid_z();
>    default: return 0;
>    }
>  }
> diff --git a/ptx-nvidiacl/lib/workitem/get_local_id.cl b/ptx-
> nvidiacl/lib/workitem/get_local_id.cl
> index f0cfdc0..f31581a 100644
> --- a/ptx-nvidiacl/lib/workitem/get_local_id.cl
> +++ b/ptx-nvidiacl/lib/workitem/get_local_id.cl
> @@ -2,9 +2,9 @@
>  
>  _CLC_DEF size_t get_local_id(uint dim) {
>    switch (dim) {
> -  case 0:  return __builtin_ptx_read_tid_x();
> -  case 1:  return __builtin_ptx_read_tid_y();
> -  case 2:  return __builtin_ptx_read_tid_z();
> +  case 0:  return __nvvm_read_ptx_sreg_tid_x();
> +  case 1:  return __nvvm_read_ptx_sreg_tid_y();
> +  case 2:  return __nvvm_read_ptx_sreg_tid_z();
>    default: return 0;
>    }
>  }
> diff --git a/ptx-nvidiacl/lib/workitem/get_local_size.cl b/ptx-
> nvidiacl/lib/workitem/get_local_size.cl
> index c3f5425..d00b0d6 100644
> --- a/ptx-nvidiacl/lib/workitem/get_local_size.cl
> +++ b/ptx-nvidiacl/lib/workitem/get_local_size.cl
> @@ -2,9 +2,9 @@
>  
>  _CLC_DEF size_t get_local_size(uint dim) {
>    switch (dim) {
> -  case 0:  return __builtin_ptx_read_ntid_x();
> -  case 1:  return __builtin_ptx_read_ntid_y();
> -  case 2:  return __builtin_ptx_read_ntid_z();
> +  case 0:  return __nvvm_read_ptx_sreg_ntid_x();
> +  case 1:  return __nvvm_read_ptx_sreg_ntid_y();
> +  case 2:  return __nvvm_read_ptx_sreg_ntid_z();
>    default: return 0;
>    }
>  }
> diff --git a/ptx-nvidiacl/lib/workitem/get_num_groups.cl b/ptx-
> nvidiacl/lib/workitem/get_num_groups.cl
> index 90bdc2e..d7abf3f 100644
> --- a/ptx-nvidiacl/lib/workitem/get_num_groups.cl
> +++ b/ptx-nvidiacl/lib/workitem/get_num_groups.cl
> @@ -2,9 +2,9 @@
>  
>  _CLC_DEF size_t get_num_groups(uint dim) {
>    switch (dim) {
> -  case 0:  return __builtin_ptx_read_nctaid_x();
> -  case 1:  return __builtin_ptx_read_nctaid_y();
> -  case 2:  return __builtin_ptx_read_nctaid_z();
> +  case 0:  return __nvvm_read_ptx_sreg_nctaid_x();
> +  case 1:  return __nvvm_read_ptx_sreg_nctaid_y();
> +  case 2:  return __nvvm_read_ptx_sreg_nctaid_z();
>    default: return 0;
>    }
>  }

ping, this is needed to fix nvptx build.

-- 
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <http://lists.llvm.org/pipermail/libclc-dev/attachments/20160721/6a120f06/attachment.sig>