1 ; RUN: llc < %s -march=ptx | FileCheck %s
3 define ptx_device i16 @tid_x() {
4 ; CHECK: mov.u16 rh0, tid.x;
6 %x = call i16 @llvm.ptx.read.tid.x()
10 define ptx_device i16 @tid_y() {
11 ; CHECK: mov.u16 rh0, tid.y;
13 %x = call i16 @llvm.ptx.read.tid.y()
17 define ptx_device i16 @tid_z() {
18 ; CHECK: mov.u16 rh0, tid.z;
20 %x = call i16 @llvm.ptx.read.tid.z()
24 define ptx_device i16 @tid_w() {
25 ; CHECK: mov.u16 rh0, tid.w;
27 %x = call i16 @llvm.ptx.read.tid.w()
31 define ptx_device void @bar_sync() {
34 call void @llvm.ptx.bar.sync(i32 0)
38 declare i16 @llvm.ptx.read.tid.x()
39 declare i16 @llvm.ptx.read.tid.y()
40 declare i16 @llvm.ptx.read.tid.z()
41 declare i16 @llvm.ptx.read.tid.w()
43 declare void @llvm.ptx.bar.sync(i32 %i)