NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx ones

The ptx spellings were removed from LLVM in r274769.

llvm-svn: 274770

1 files changed, 13 insertions, 13 deletions

diff --git a/clang/lib/Headers/cuda_builtin_vars.h b/clang/lib/Headers/cuda_builtin_vars.h
index ec830864036..6f5eb9c78d8 100644
--- a/clang/lib/Headers/cuda_builtin_vars.h
+++ b/clang/lib/Headers/cuda_builtin_vars.h
@@ -37,7 +37,7 @@ struct dim3;
 // Example:
 //    int x = threadIdx.x;
 // IR output:
-//  %0 = call i32 @llvm.ptx.read.tid.x() #3
+//  %0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #3
 // PTX output:
 //  mov.u32     %r2, %tid.x;
 
@@ -64,9 +64,9 @@ struct dim3;
   __attribute__((device)) TypeName *operator&() const __DELETE
 
 struct __cuda_builtin_threadIdx_t {
-  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_tid_x());
-  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_tid_y());
-  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_tid_z());
+  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_tid_x());
+  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_tid_y());
+  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_tid_z());
   // threadIdx should be convertible to uint3 (in fact in nvcc, it *is* a
   // uint3).  This function is defined after we pull in vector_types.h.
   __attribute__((device)) operator uint3() const;
@@ -75,9 +75,9 @@ private:
 };
 
 struct __cuda_builtin_blockIdx_t {
-  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_ctaid_x());
-  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_ctaid_y());
-  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_ctaid_z());
+  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ctaid_x());
+  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ctaid_y());
+  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ctaid_z());
   // blockIdx should be convertible to uint3 (in fact in nvcc, it *is* a
   // uint3).  This function is defined after we pull in vector_types.h.
   __attribute__((device)) operator uint3() const;
@@ -86,9 +86,9 @@ private:
 };
 
 struct __cuda_builtin_blockDim_t {
-  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_ntid_x());
-  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_ntid_y());
-  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_ntid_z());
+  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ntid_x());
+  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ntid_y());
+  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ntid_z());
   // blockDim should be convertible to dim3 (in fact in nvcc, it *is* a
   // dim3).  This function is defined after we pull in vector_types.h.
   __attribute__((device)) operator dim3() const;
@@ -97,9 +97,9 @@ private:
 };
 
 struct __cuda_builtin_gridDim_t {
-  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_nctaid_x());
-  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_nctaid_y());
-  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_nctaid_z());
+  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_nctaid_x());
+  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_nctaid_y());
+  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_nctaid_z());
   // gridDim should be convertible to dim3 (in fact in nvcc, it *is* a
   // dim3).  This function is defined after we pull in vector_types.h.
   __attribute__((device)) operator dim3() const;