[libclc] 36/58: ptx: Fix builtin names after clang r274770

[Andreas Boll]

This is an automated email from the git hooks/post-receive script.

aboll-guest pushed a commit to branch master
in repository libclc.

commit 3d39fb557be2e349ded2d2055211b2757627f40b
Author: Jan Vesely <jan.vesely at rutgers.edu>
Date:   Fri Jul 22 15:00:08 2016 +0000

    ptx: Fix builtin names after clang r274770
    
    Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
    Acked-By: Aaron Watry <awatry at gmail.com>
    
    git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@276423 91177308-0d34-0410-b5e6-96231b3b80d8
---
 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;
   }
 }

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/libclc.git