diff options
| author | Justin Bogner <mail@justinbogner.com> | 2016-07-07 16:41:08 +0000 |
|---|---|---|
| committer | Justin Bogner <mail@justinbogner.com> | 2016-07-07 16:41:08 +0000 |
| commit | 29dc237122635f074aff00cde059684fc521c73d (patch) | |
| tree | 42ae2d3fb5d649711375a95a984fe34374795b33 /test/CodeGenCUDA | |
| parent | 236751c31c3a8914b3994645ec579aef09e29572 (diff) | |
NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx ones
The ptx spellings were removed from LLVM in r274769.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@274770 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
| -rw-r--r-- | test/CodeGenCUDA/cuda-builtin-vars.cu | 24 |
1 files changed, 12 insertions, 12 deletions
diff --git a/test/CodeGenCUDA/cuda-builtin-vars.cu b/test/CodeGenCUDA/cuda-builtin-vars.cu index 834e16d04d..c2159f5af1 100644 --- a/test/CodeGenCUDA/cuda-builtin-vars.cu +++ b/test/CodeGenCUDA/cuda-builtin-vars.cu @@ -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, |