diff options
Diffstat (limited to 'test/CodeGen/PTX/intrinsic.ll')
-rw-r--r-- | test/CodeGen/PTX/intrinsic.ll | 96 |
1 files changed, 96 insertions, 0 deletions
diff --git a/test/CodeGen/PTX/intrinsic.ll b/test/CodeGen/PTX/intrinsic.ll index 24c328b773..804d8b5f95 100644 --- a/test/CodeGen/PTX/intrinsic.ll +++ b/test/CodeGen/PTX/intrinsic.ll @@ -28,6 +28,90 @@ define ptx_device i16 @tid_w() { ret i16 %x } +define ptx_device i16 @ntid_x() { +; CHECK: mov.u16 rh0, ntid.x; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ntid.x() + ret i16 %x +} + +define ptx_device i16 @ntid_y() { +; CHECK: mov.u16 rh0, ntid.y; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ntid.y() + ret i16 %x +} + +define ptx_device i16 @ntid_z() { +; CHECK: mov.u16 rh0, ntid.z; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ntid.z() + ret i16 %x +} + +define ptx_device i16 @ntid_w() { +; CHECK: mov.u16 rh0, ntid.w; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ntid.w() + ret i16 %x +} + +define ptx_device i16 @ctaid_x() { +; CHECK: mov.u16 rh0, ctaid.x; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ctaid.x() + ret i16 %x +} + +define ptx_device i16 @ctaid_y() { +; CHECK: mov.u16 rh0, ctaid.y; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ctaid.y() + ret i16 %x +} + +define ptx_device i16 @ctaid_z() { +; CHECK: mov.u16 rh0, ctaid.z; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ctaid.z() + ret i16 %x +} + +define ptx_device i16 @ctaid_w() { +; CHECK: mov.u16 rh0, ctaid.w; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.ctaid.w() + ret i16 %x +} + +define ptx_device i16 @nctaid_x() { +; CHECK: mov.u16 rh0, nctaid.x; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.nctaid.x() + ret i16 %x +} + +define ptx_device i16 @nctaid_y() { +; CHECK: mov.u16 rh0, nctaid.y; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.nctaid.y() + ret i16 %x +} + +define ptx_device i16 @nctaid_z() { +; CHECK: mov.u16 rh0, nctaid.z; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.nctaid.z() + ret i16 %x +} + +define ptx_device i16 @nctaid_w() { +; CHECK: mov.u16 rh0, nctaid.w; +; CHECK-NEXT: ret; + %x = call i16 @llvm.ptx.read.nctaid.w() + ret i16 %x +} + define ptx_device void @bar_sync() { ; CHECK: bar.sync 0 ; CHECK-NEXT: ret; @@ -39,5 +123,17 @@ declare i16 @llvm.ptx.read.tid.x() declare i16 @llvm.ptx.read.tid.y() declare i16 @llvm.ptx.read.tid.z() declare i16 @llvm.ptx.read.tid.w() +declare i16 @llvm.ptx.read.ntid.x() +declare i16 @llvm.ptx.read.ntid.y() +declare i16 @llvm.ptx.read.ntid.z() +declare i16 @llvm.ptx.read.ntid.w() +declare i16 @llvm.ptx.read.ctaid.x() +declare i16 @llvm.ptx.read.ctaid.y() +declare i16 @llvm.ptx.read.ctaid.z() +declare i16 @llvm.ptx.read.ctaid.w() +declare i16 @llvm.ptx.read.nctaid.x() +declare i16 @llvm.ptx.read.nctaid.y() +declare i16 @llvm.ptx.read.nctaid.z() +declare i16 @llvm.ptx.read.nctaid.w() declare void @llvm.ptx.bar.sync(i32 %i) |