diff options
Diffstat (limited to 'llvm/test/CodeGen/NVPTX')
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/intrinsic-old.ll | 67 |
1 files changed, 54 insertions, 13 deletions
diff --git a/llvm/test/CodeGen/NVPTX/intrinsic-old.ll b/llvm/test/CodeGen/NVPTX/intrinsic-old.ll index 3c51776c0ec..5c73f44d075 100644 --- a/llvm/test/CodeGen/NVPTX/intrinsic-old.ll +++ b/llvm/test/CodeGen/NVPTX/intrinsic-old.ll @@ -1,8 +1,14 @@ ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s +; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda -nvvm-intr-range \ +; RUN: | FileCheck --check-prefix=RANGE --check-prefix=RANGE_20 %s +; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda \ +; RUN: -nvvm-intr-range -nvvm-intr-range-sm=30 \ +; RUN: | FileCheck --check-prefix=RANGE --check-prefix=RANGE_30 %s define ptx_device i32 @test_tid_x() { ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x; +; RANGE: call i32 @llvm.ptx.read.tid.x(), !range ![[BLK_IDX_XY:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.tid.x() ret i32 %x @@ -10,6 +16,7 @@ define ptx_device i32 @test_tid_x() { define ptx_device i32 @test_tid_y() { ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.y; +; RANGE: call i32 @llvm.ptx.read.tid.y(), !range ![[BLK_IDX_XY]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.tid.y() ret i32 %x @@ -17,6 +24,7 @@ define ptx_device i32 @test_tid_y() { define ptx_device i32 @test_tid_z() { ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.z; +; RANGE: call i32 @llvm.ptx.read.tid.z(), !range ![[BLK_IDX_Z:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.tid.z() ret i32 %x @@ -31,6 +39,7 @@ define ptx_device i32 @test_tid_w() { define ptx_device i32 @test_ntid_x() { ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.x; +; RANGE: call i32 @llvm.ptx.read.ntid.x(), !range ![[BLK_SIZE_XY:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.ntid.x() ret i32 %x @@ -38,6 +47,7 @@ define ptx_device i32 @test_ntid_x() { define ptx_device i32 @test_ntid_y() { ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.y; +; RANGE: call i32 @llvm.ptx.read.ntid.y(), !range ![[BLK_SIZE_XY]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.ntid.y() ret i32 %x @@ -45,6 +55,7 @@ define ptx_device i32 @test_ntid_y() { define ptx_device i32 @test_ntid_z() { ; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.z; +; RANGE: call i32 @llvm.ptx.read.ntid.z(), !range ![[BLK_SIZE_Z:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.ntid.z() ret i32 %x @@ -59,11 +70,20 @@ define ptx_device i32 @test_ntid_w() { define ptx_device i32 @test_laneid() { ; CHECK: mov.u32 %r{{[0-9]+}}, %laneid; +; RANGE: call i32 @llvm.ptx.read.laneid(), !range ![[LANEID:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.laneid() ret i32 %x } +define ptx_device i32 @test_warpsize() { +; CHECK: mov.u32 %r{{[0-9]+}}, WARP_SZ; +; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.warpsize(), !range ![[WARPSIZE:[0-9]+]] +; CHECK: ret; + %x = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() + ret i32 %x +} + define ptx_device i32 @test_warpid() { ; CHECK: mov.u32 %r{{[0-9]+}}, %warpid; ; CHECK: ret; @@ -78,15 +98,9 @@ define ptx_device i32 @test_nwarpid() { ret i32 %x } -define ptx_device i32 @test_ctaid_x() { -; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x; -; CHECK: ret; - %x = call i32 @llvm.ptx.read.ctaid.x() - ret i32 %x -} - define ptx_device i32 @test_ctaid_y() { ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.y; +; RANGE: call i32 @llvm.ptx.read.ctaid.y(), !range ![[GRID_IDX_YZ:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.ctaid.y() ret i32 %x @@ -94,27 +108,31 @@ define ptx_device i32 @test_ctaid_y() { define ptx_device i32 @test_ctaid_z() { ; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.z; +; RANGE: call i32 @llvm.ptx.read.ctaid.z(), !range ![[GRID_IDX_YZ]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.ctaid.z() ret i32 %x } -define ptx_device i32 @test_ctaid_w() { -; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w; +define ptx_device i32 @test_ctaid_x() { +; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x; +; RANGE_30: call i32 @llvm.ptx.read.ctaid.x(), !range ![[GRID_IDX_X:[0-9]+]] +; RANGE_20: call i32 @llvm.ptx.read.ctaid.x(), !range ![[GRID_IDX_YZ]] ; CHECK: ret; - %x = call i32 @llvm.ptx.read.ctaid.w() + %x = call i32 @llvm.ptx.read.ctaid.x() ret i32 %x } -define ptx_device i32 @test_nctaid_x() { -; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x; +define ptx_device i32 @test_ctaid_w() { +; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w; ; CHECK: ret; - %x = call i32 @llvm.ptx.read.nctaid.x() + %x = call i32 @llvm.ptx.read.ctaid.w() ret i32 %x } define ptx_device i32 @test_nctaid_y() { ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.y; +; RANGE: call i32 @llvm.ptx.read.nctaid.y(), !range ![[GRID_SIZE_YZ:[0-9]+]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.nctaid.y() ret i32 %x @@ -122,11 +140,22 @@ define ptx_device i32 @test_nctaid_y() { define ptx_device i32 @test_nctaid_z() { ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.z; +; RANGE: call i32 @llvm.ptx.read.nctaid.z(), !range ![[GRID_SIZE_YZ]] ; CHECK: ret; %x = call i32 @llvm.ptx.read.nctaid.z() ret i32 %x } +define ptx_device i32 @test_nctaid_x() { +; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x; +; RANGE_30: call i32 @llvm.ptx.read.nctaid.x(), !range ![[GRID_SIZE_X:[0-9]+]] +; RANGE_20: call i32 @llvm.ptx.read.nctaid.x(), !range ![[GRID_SIZE_YZ]] +; CHECK: ret; + %x = call i32 @llvm.ptx.read.nctaid.x() + ret i32 %x +} + + define ptx_device i32 @test_nctaid_w() { ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w; ; CHECK: ret; @@ -248,6 +277,7 @@ declare i32 @llvm.ptx.read.ntid.y() declare i32 @llvm.ptx.read.ntid.z() declare i32 @llvm.ptx.read.ntid.w() +declare i32 @llvm.nvvm.read.ptx.sreg.warpsize() declare i32 @llvm.ptx.read.laneid() declare i32 @llvm.ptx.read.warpid() declare i32 @llvm.ptx.read.nwarpid() @@ -280,3 +310,14 @@ declare i32 @llvm.ptx.read.pm2() declare i32 @llvm.ptx.read.pm3() declare void @llvm.ptx.bar.sync(i32 %i) + +; RANGE-DAG: ![[BLK_IDX_XY]] = !{i32 0, i32 1024} +; RANGE-DAG: ![[BLK_IDX_Z]] = !{i32 0, i32 64} +; RANGE-DAG: ![[BLK_SIZE_XY]] = !{i32 1, i32 1025} +; RANGE-DAG: ![[BLK_SIZE_Z]] = !{i32 1, i32 65} +; RANGE-DAG: ![[LANEID]] = !{i32 0, i32 32} +; RANGE-DAG: ![[WARPSIZE]] = !{i32 32, i32 33} +; RANGE_30-DAG: ![[GRID_IDX_X]] = !{i32 0, i32 2147483647} +; RANGE-DAG: ![[GRID_IDX_YZ]] = !{i32 0, i32 65535} +; RANGE_30-DAG: ![[GRID_SIZE_X]] = !{i32 1, i32 -2147483648} +; RANGE-DAG: ![[GRID_SIZE_YZ]] = !{i32 1, i32 65536} |
