[libclc] r276423 - ptx: Fix builtin names after clang r274770
Jan Vesely via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 22 08:00:10 PDT 2016
Author: jvesely
Date: Fri Jul 22 10:00:08 2016
New Revision: 276423
URL: http://llvm.org/viewvc/llvm-project?rev=276423&view=rev
Log:
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>
Modified:
libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl
libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl
libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl
libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl
libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl
Modified: libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl?rev=276423&r1=276422&r2=276423&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl (original)
+++ libclc/trunk/ptx-nvidiacl/lib/synchronization/barrier.cl Fri Jul 22 10:00:08 2016
@@ -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();
}
}
Modified: libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl?rev=276423&r1=276422&r2=276423&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_group_id.cl Fri Jul 22 10:00:08 2016
@@ -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;
}
}
Modified: libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl?rev=276423&r1=276422&r2=276423&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_id.cl Fri Jul 22 10:00:08 2016
@@ -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;
}
}
Modified: libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl?rev=276423&r1=276422&r2=276423&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_local_size.cl Fri Jul 22 10:00:08 2016
@@ -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;
}
}
Modified: libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl?rev=276423&r1=276422&r2=276423&view=diff
==============================================================================
--- libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl (original)
+++ libclc/trunk/ptx-nvidiacl/lib/workitem/get_num_groups.cl Fri Jul 22 10:00:08 2016
@@ -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;
}
}
More information about the cfe-commits
mailing list