r274770 - NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx ones
Justin Bogner via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 7 09:41:08 PDT 2016
Author: bogner
Date: Thu Jul 7 11:41:08 2016
New Revision: 274770
URL: http://llvm.org/viewvc/llvm-project?rev=274770&view=rev
Log:
NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx ones
The ptx spellings were removed from LLVM in r274769.
Modified:
cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
cfe/trunk/lib/Headers/cuda_builtin_vars.h
cfe/trunk/test/CodeGen/builtins-nvptx.c
cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu
cfe/trunk/test/SemaCUDA/builtins.cu
Modified: cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def?rev=274770&r1=274769&r2=274770&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def Thu Jul 7 11:41:08 2016
@@ -14,53 +14,50 @@
// The format of this database matches clang/Basic/Builtins.def.
-// Builtins retained from previous PTX back-end
-BUILTIN(__builtin_ptx_read_tid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_tid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_tid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_tid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_ntid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_ntid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_ntid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_ntid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_ctaid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_ctaid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_ctaid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_ctaid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_nctaid_x, "i", "nc")
-BUILTIN(__builtin_ptx_read_nctaid_y, "i", "nc")
-BUILTIN(__builtin_ptx_read_nctaid_z, "i", "nc")
-BUILTIN(__builtin_ptx_read_nctaid_w, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_laneid, "i", "nc")
-BUILTIN(__builtin_ptx_read_warpid, "i", "nc")
-BUILTIN(__builtin_ptx_read_nwarpid, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_smid, "i", "nc")
-BUILTIN(__builtin_ptx_read_nsmid, "i", "nc")
-BUILTIN(__builtin_ptx_read_gridid, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_lanemask_eq, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_le, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_lt, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_ge, "i", "nc")
-BUILTIN(__builtin_ptx_read_lanemask_gt, "i", "nc")
-
-BUILTIN(__builtin_ptx_read_clock, "i", "n")
-BUILTIN(__builtin_ptx_read_clock64, "LLi", "n")
-
-BUILTIN(__builtin_ptx_read_pm0, "i", "n")
-BUILTIN(__builtin_ptx_read_pm1, "i", "n")
-BUILTIN(__builtin_ptx_read_pm2, "i", "n")
-BUILTIN(__builtin_ptx_read_pm3, "i", "n")
+// Special Registers
-BUILTIN(__builtin_ptx_bar_sync, "vi", "n")
+BUILTIN(__nvvm_read_ptx_sreg_tid_x, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_tid_y, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_tid_z, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_tid_w, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_ntid_x, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_ntid_y, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_ntid_z, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_ntid_w, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_ctaid_x, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_ctaid_y, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_ctaid_z, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_ctaid_w, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_nctaid_x, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_nctaid_y, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_nctaid_z, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_nctaid_w, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_laneid, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_warpid, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_nwarpid, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_smid, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_nsmid, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_gridid, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_lanemask_eq, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_lanemask_le, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_lanemask_lt, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_lanemask_ge, "i", "nc")
+BUILTIN(__nvvm_read_ptx_sreg_lanemask_gt, "i", "nc")
+
+BUILTIN(__nvvm_read_ptx_sreg_clock, "i", "n")
+BUILTIN(__nvvm_read_ptx_sreg_clock64, "LLi", "n")
+
+BUILTIN(__nvvm_read_ptx_sreg_pm0, "i", "n")
+BUILTIN(__nvvm_read_ptx_sreg_pm1, "i", "n")
+BUILTIN(__nvvm_read_ptx_sreg_pm2, "i", "n")
+BUILTIN(__nvvm_read_ptx_sreg_pm3, "i", "n")
-
-// Builtins exposed as part of NVVM
// MISC
BUILTIN(__nvvm_clz_i, "ii", "")
@@ -396,11 +393,11 @@ BUILTIN(__nvvm_bitcast_d2ll, "LLid", "")
// Sync
-BUILTIN(__syncthreads, "v", "")
BUILTIN(__nvvm_bar0, "v", "")
BUILTIN(__nvvm_bar0_popc, "ii", "")
BUILTIN(__nvvm_bar0_and, "ii", "")
BUILTIN(__nvvm_bar0_or, "ii", "")
+BUILTIN(__nvvm_bar_sync, "vi", "n")
// Shuffle
Modified: cfe/trunk/lib/Headers/cuda_builtin_vars.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/cuda_builtin_vars.h?rev=274770&r1=274769&r2=274770&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/cuda_builtin_vars.h (original)
+++ cfe/trunk/lib/Headers/cuda_builtin_vars.h Thu Jul 7 11:41:08 2016
@@ -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;
Modified: cfe/trunk/test/CodeGen/builtins-nvptx.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx.c?rev=274770&r1=274769&r2=274770&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/builtins-nvptx.c (original)
+++ cfe/trunk/test/CodeGen/builtins-nvptx.c Thu Jul 7 11:41:08 2016
@@ -11,15 +11,15 @@
__device__ int read_tid() {
-// CHECK: call i32 @llvm.ptx.read.tid.x()
-// CHECK: call i32 @llvm.ptx.read.tid.y()
-// CHECK: call i32 @llvm.ptx.read.tid.z()
-// CHECK: call i32 @llvm.ptx.read.tid.w()
-
- int x = __builtin_ptx_read_tid_x();
- int y = __builtin_ptx_read_tid_y();
- int z = __builtin_ptx_read_tid_z();
- int w = __builtin_ptx_read_tid_w();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.w()
+
+ int x = __nvvm_read_ptx_sreg_tid_x();
+ int y = __nvvm_read_ptx_sreg_tid_y();
+ int z = __nvvm_read_ptx_sreg_tid_z();
+ int w = __nvvm_read_ptx_sreg_tid_w();
return x + y + z + w;
@@ -27,15 +27,15 @@ __device__ int read_tid() {
__device__ int read_ntid() {
-// CHECK: call i32 @llvm.ptx.read.ntid.x()
-// CHECK: call i32 @llvm.ptx.read.ntid.y()
-// CHECK: call i32 @llvm.ptx.read.ntid.z()
-// CHECK: call i32 @llvm.ptx.read.ntid.w()
-
- int x = __builtin_ptx_read_ntid_x();
- int y = __builtin_ptx_read_ntid_y();
- int z = __builtin_ptx_read_ntid_z();
- int w = __builtin_ptx_read_ntid_w();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.w()
+
+ int x = __nvvm_read_ptx_sreg_ntid_x();
+ int y = __nvvm_read_ptx_sreg_ntid_y();
+ int z = __nvvm_read_ptx_sreg_ntid_z();
+ int w = __nvvm_read_ptx_sreg_ntid_w();
return x + y + z + w;
@@ -43,15 +43,15 @@ __device__ int read_ntid() {
__device__ int read_ctaid() {
-// CHECK: call i32 @llvm.ptx.read.ctaid.x()
-// CHECK: call i32 @llvm.ptx.read.ctaid.y()
-// CHECK: call i32 @llvm.ptx.read.ctaid.z()
-// CHECK: call i32 @llvm.ptx.read.ctaid.w()
-
- int x = __builtin_ptx_read_ctaid_x();
- int y = __builtin_ptx_read_ctaid_y();
- int z = __builtin_ptx_read_ctaid_z();
- int w = __builtin_ptx_read_ctaid_w();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.w()
+
+ int x = __nvvm_read_ptx_sreg_ctaid_x();
+ int y = __nvvm_read_ptx_sreg_ctaid_y();
+ int z = __nvvm_read_ptx_sreg_ctaid_z();
+ int w = __nvvm_read_ptx_sreg_ctaid_w();
return x + y + z + w;
@@ -59,15 +59,15 @@ __device__ int read_ctaid() {
__device__ int read_nctaid() {
-// CHECK: call i32 @llvm.ptx.read.nctaid.x()
-// CHECK: call i32 @llvm.ptx.read.nctaid.y()
-// CHECK: call i32 @llvm.ptx.read.nctaid.z()
-// CHECK: call i32 @llvm.ptx.read.nctaid.w()
-
- int x = __builtin_ptx_read_nctaid_x();
- int y = __builtin_ptx_read_nctaid_y();
- int z = __builtin_ptx_read_nctaid_z();
- int w = __builtin_ptx_read_nctaid_w();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.w()
+
+ int x = __nvvm_read_ptx_sreg_nctaid_x();
+ int y = __nvvm_read_ptx_sreg_nctaid_y();
+ int z = __nvvm_read_ptx_sreg_nctaid_z();
+ int w = __nvvm_read_ptx_sreg_nctaid_w();
return x + y + z + w;
@@ -75,19 +75,19 @@ __device__ int read_nctaid() {
__device__ int read_ids() {
-// CHECK: call i32 @llvm.ptx.read.laneid()
-// CHECK: call i32 @llvm.ptx.read.warpid()
-// CHECK: call i32 @llvm.ptx.read.nwarpid()
-// CHECK: call i32 @llvm.ptx.read.smid()
-// CHECK: call i32 @llvm.ptx.read.nsmid()
-// CHECK: call i32 @llvm.ptx.read.gridid()
-
- int a = __builtin_ptx_read_laneid();
- int b = __builtin_ptx_read_warpid();
- int c = __builtin_ptx_read_nwarpid();
- int d = __builtin_ptx_read_smid();
- int e = __builtin_ptx_read_nsmid();
- int f = __builtin_ptx_read_gridid();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.laneid()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpid()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nwarpid()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.smid()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nsmid()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.gridid()
+
+ int a = __nvvm_read_ptx_sreg_laneid();
+ int b = __nvvm_read_ptx_sreg_warpid();
+ int c = __nvvm_read_ptx_sreg_nwarpid();
+ int d = __nvvm_read_ptx_sreg_smid();
+ int e = __nvvm_read_ptx_sreg_nsmid();
+ int f = __nvvm_read_ptx_sreg_gridid();
return a + b + c + d + e + f;
@@ -95,17 +95,17 @@ __device__ int read_ids() {
__device__ int read_lanemasks() {
-// CHECK: call i32 @llvm.ptx.read.lanemask.eq()
-// CHECK: call i32 @llvm.ptx.read.lanemask.le()
-// CHECK: call i32 @llvm.ptx.read.lanemask.lt()
-// CHECK: call i32 @llvm.ptx.read.lanemask.ge()
-// CHECK: call i32 @llvm.ptx.read.lanemask.gt()
-
- int a = __builtin_ptx_read_lanemask_eq();
- int b = __builtin_ptx_read_lanemask_le();
- int c = __builtin_ptx_read_lanemask_lt();
- int d = __builtin_ptx_read_lanemask_ge();
- int e = __builtin_ptx_read_lanemask_gt();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.eq()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.le()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.lt()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.ge()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.gt()
+
+ int a = __nvvm_read_ptx_sreg_lanemask_eq();
+ int b = __nvvm_read_ptx_sreg_lanemask_le();
+ int c = __nvvm_read_ptx_sreg_lanemask_lt();
+ int d = __nvvm_read_ptx_sreg_lanemask_ge();
+ int e = __nvvm_read_ptx_sreg_lanemask_gt();
return a + b + c + d + e;
@@ -113,26 +113,26 @@ __device__ int read_lanemasks() {
__device__ long long read_clocks() {
-// CHECK: call i32 @llvm.ptx.read.clock()
-// CHECK: call i64 @llvm.ptx.read.clock64()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clock()
+// CHECK: call i64 @llvm.nvvm.read.ptx.sreg.clock64()
- int a = __builtin_ptx_read_clock();
- long long b = __builtin_ptx_read_clock64();
+ int a = __nvvm_read_ptx_sreg_clock();
+ long long b = __nvvm_read_ptx_sreg_clock64();
return a + b;
}
__device__ int read_pms() {
-// CHECK: call i32 @llvm.ptx.read.pm0()
-// CHECK: call i32 @llvm.ptx.read.pm1()
-// CHECK: call i32 @llvm.ptx.read.pm2()
-// CHECK: call i32 @llvm.ptx.read.pm3()
-
- int a = __builtin_ptx_read_pm0();
- int b = __builtin_ptx_read_pm1();
- int c = __builtin_ptx_read_pm2();
- int d = __builtin_ptx_read_pm3();
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm0()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm1()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm2()
+// CHECK: call i32 @llvm.nvvm.read.ptx.sreg.pm3()
+
+ int a = __nvvm_read_ptx_sreg_pm0();
+ int b = __nvvm_read_ptx_sreg_pm1();
+ int c = __nvvm_read_ptx_sreg_pm2();
+ int d = __nvvm_read_ptx_sreg_pm3();
return a + b + c + d;
@@ -140,9 +140,9 @@ __device__ int read_pms() {
__device__ void sync() {
-// CHECK: call void @llvm.ptx.bar.sync(i32 0)
+// CHECK: call void @llvm.nvvm.bar.sync(i32 0)
- __builtin_ptx_bar_sync(0);
+ __nvvm_bar_sync(0);
}
Modified: cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu?rev=274770&r1=274769&r2=274770&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu Thu Jul 7 11:41:08 2016
@@ -6,21 +6,21 @@
__attribute__((global))
void kernel(int *out) {
int i = 0;
- out[i++] = threadIdx.x; // CHECK: call i32 @llvm.ptx.read.tid.x()
- out[i++] = threadIdx.y; // CHECK: call i32 @llvm.ptx.read.tid.y()
- out[i++] = threadIdx.z; // CHECK: call i32 @llvm.ptx.read.tid.z()
+ out[i++] = threadIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+ out[i++] = threadIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
+ out[i++] = threadIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
- out[i++] = blockIdx.x; // CHECK: call i32 @llvm.ptx.read.ctaid.x()
- out[i++] = blockIdx.y; // CHECK: call i32 @llvm.ptx.read.ctaid.y()
- out[i++] = blockIdx.z; // CHECK: call i32 @llvm.ptx.read.ctaid.z()
+ out[i++] = blockIdx.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
+ out[i++] = blockIdx.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
+ out[i++] = blockIdx.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
- out[i++] = blockDim.x; // CHECK: call i32 @llvm.ptx.read.ntid.x()
- out[i++] = blockDim.y; // CHECK: call i32 @llvm.ptx.read.ntid.y()
- out[i++] = blockDim.z; // CHECK: call i32 @llvm.ptx.read.ntid.z()
+ out[i++] = blockDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
+ out[i++] = blockDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
+ out[i++] = blockDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
- out[i++] = gridDim.x; // CHECK: call i32 @llvm.ptx.read.nctaid.x()
- out[i++] = gridDim.y; // CHECK: call i32 @llvm.ptx.read.nctaid.y()
- out[i++] = gridDim.z; // CHECK: call i32 @llvm.ptx.read.nctaid.z()
+ out[i++] = gridDim.x; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
+ out[i++] = gridDim.y; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
+ out[i++] = gridDim.z; // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
out[i++] = warpSize; // CHECK: store i32 32,
Modified: cfe/trunk/test/SemaCUDA/builtins.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/builtins.cu?rev=274770&r1=274769&r2=274770&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/builtins.cu (original)
+++ cfe/trunk/test/SemaCUDA/builtins.cu Thu Jul 7 11:41:08 2016
@@ -18,13 +18,13 @@
void hf() {
int x = __builtin_ia32_rdtsc();
- int y = __builtin_ptx_read_tid_x(); // expected-note {{'__builtin_ptx_read_tid_x' declared here}}
- // expected-error at -1 {{reference to __device__ function '__builtin_ptx_read_tid_x' in __host__ function}}
+ int y = __nvvm_read_ptx_sreg_tid_x(); // expected-note {{'__nvvm_read_ptx_sreg_tid_x' declared here}}
+ // expected-error at -1 {{reference to __device__ function '__nvvm_read_ptx_sreg_tid_x' in __host__ function}}
x = __builtin_abs(1);
}
__attribute__((device)) void df() {
- int x = __builtin_ptx_read_tid_x();
+ int x = __nvvm_read_ptx_sreg_tid_x();
int y = __builtin_ia32_rdtsc(); // expected-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
// expected-note at 20 {{'__builtin_ia32_rdtsc' declared here}}
x = __builtin_abs(1);
More information about the cfe-commits
mailing list