diff options
author | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-22 15:00:08 +0000 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2016-07-22 15:00:08 +0000 |
commit | 3d39fb557be2e349ded2d2055211b2757627f40b (patch) | |
tree | 10aa445bda8a7f4f220c35e11e44bcb37a69d947 | |
parent | a3d6cec243bbe0446e8500cfd6a241c3357aca41 (diff) |
ptx: Fix builtin names after clang r274770
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-By: Aaron Watry <awatry@gmail.com>
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@276423 91177308-0d34-0410-b5e6-96231b3b80d8
-rw-r--r-- | ptx-nvidiacl/lib/synchronization/barrier.cl | 2 | ||||
-rw-r--r-- | ptx-nvidiacl/lib/workitem/get_group_id.cl | 6 | ||||
-rw-r--r-- | ptx-nvidiacl/lib/workitem/get_local_id.cl | 6 | ||||
-rw-r--r-- | ptx-nvidiacl/lib/workitem/get_local_size.cl | 6 | ||||
-rw-r--r-- | 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; } } |