303 lines · plain
1; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck -allow-deprecated-dag-overlap %s2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck -allow-deprecated-dag-overlap %s3; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}5 6define ptx_device i32 @test_tid_x() {7; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x;8; CHECK: ret;9 %x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()10 ret i32 %x11}12 13define ptx_device i32 @test_tid_y() {14; CHECK: mov.u32 %r{{[0-9]+}}, %tid.y;15; CHECK: ret;16 %x = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()17 ret i32 %x18}19 20define ptx_device i32 @test_tid_z() {21; CHECK: mov.u32 %r{{[0-9]+}}, %tid.z;22; CHECK: ret;23 %x = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()24 ret i32 %x25}26 27define ptx_device i32 @test_tid_w() {28; CHECK: mov.u32 %r{{[0-9]+}}, %tid.w;29; CHECK: ret;30 %x = call i32 @llvm.nvvm.read.ptx.sreg.tid.w()31 ret i32 %x32}33 34define ptx_device i32 @test_ntid_x() {35; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.x;36; CHECK: ret;37 %x = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()38 ret i32 %x39}40 41define ptx_device i32 @test_ntid_y() {42; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.y;43; CHECK: ret;44 %x = call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()45 ret i32 %x46}47 48define ptx_device i32 @test_ntid_z() {49; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.z;50; CHECK: ret;51 %x = call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()52 ret i32 %x53}54 55define ptx_device i32 @test_ntid_w() {56; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.w;57; CHECK: ret;58 %x = call i32 @llvm.nvvm.read.ptx.sreg.ntid.w()59 ret i32 %x60}61 62define ptx_device i32 @test_laneid() {63; CHECK: mov.u32 %r{{[0-9]+}}, %laneid;64; CHECK: ret;65 %x = call i32 @llvm.nvvm.read.ptx.sreg.laneid()66 ret i32 %x67}68 69define ptx_device i32 @test_warpsize() {70; CHECK: mov.u32 %r{{[0-9]+}}, WARP_SZ;71; CHECK: ret;72 %x = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()73 ret i32 %x74}75 76define ptx_device i32 @test_warpid() {77; CHECK: mov.u32 %r{{[0-9]+}}, %warpid;78; CHECK: ret;79 %x = call i32 @llvm.nvvm.read.ptx.sreg.warpid()80 ret i32 %x81}82 83define ptx_device i32 @test_nwarpid() {84; CHECK: mov.u32 %r{{[0-9]+}}, %nwarpid;85; CHECK: ret;86 %x = call i32 @llvm.nvvm.read.ptx.sreg.nwarpid()87 ret i32 %x88}89 90define ptx_device i32 @test_ctaid_y() {91; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.y;92; CHECK: ret;93 %x = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()94 ret i32 %x95}96 97define ptx_device i32 @test_ctaid_z() {98; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.z;99; CHECK: ret;100 %x = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()101 ret i32 %x102}103 104define ptx_device i32 @test_ctaid_x() {105; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x;106; CHECK: ret;107 %x = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()108 ret i32 %x109}110 111define ptx_device i32 @test_ctaid_w() {112; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w;113; CHECK: ret;114 %x = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.w()115 ret i32 %x116}117 118define ptx_device i32 @test_nctaid_y() {119; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.y;120; CHECK: ret;121 %x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()122 ret i32 %x123}124 125define ptx_device i32 @test_nctaid_z() {126; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.z;127; CHECK: ret;128 %x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()129 ret i32 %x130}131 132define ptx_device i32 @test_nctaid_x() {133; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;134; CHECK: ret;135 %x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()136 ret i32 %x137}138 139define ptx_device i32 @test_already_has_range_md() {140; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;141; CHECK: ret;142 %x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x(), !range !0143 ret i32 %x144}145 146 147define ptx_device i32 @test_nctaid_w() {148; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w;149; CHECK: ret;150 %x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.w()151 ret i32 %x152}153 154define ptx_device i32 @test_smid() {155; CHECK: mov.u32 %r{{[0-9]+}}, %smid;156; CHECK: ret;157 %x = call i32 @llvm.nvvm.read.ptx.sreg.smid()158 ret i32 %x159}160 161define ptx_device i32 @test_nsmid() {162; CHECK: mov.u32 %r{{[0-9]+}}, %nsmid;163; CHECK: ret;164 %x = call i32 @llvm.nvvm.read.ptx.sreg.nsmid()165 ret i32 %x166}167 168define ptx_device i32 @test_gridid() {169; CHECK: mov.u32 %r{{[0-9]+}}, %gridid;170; CHECK: ret;171 %x = call i32 @llvm.nvvm.read.ptx.sreg.gridid()172 ret i32 %x173}174 175define ptx_device i32 @test_lanemask_eq() {176; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_eq;177; CHECK: ret;178 %x = call i32 @llvm.nvvm.read.ptx.sreg.lanemask.eq()179 ret i32 %x180}181 182define ptx_device i32 @test_lanemask_le() {183; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_le;184; CHECK: ret;185 %x = call i32 @llvm.nvvm.read.ptx.sreg.lanemask.le()186 ret i32 %x187}188 189define ptx_device i32 @test_lanemask_lt() {190; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_lt;191; CHECK: ret;192 %x = call i32 @llvm.nvvm.read.ptx.sreg.lanemask.lt()193 ret i32 %x194}195 196define ptx_device i32 @test_lanemask_ge() {197; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_ge;198; CHECK: ret;199 %x = call i32 @llvm.nvvm.read.ptx.sreg.lanemask.ge()200 ret i32 %x201}202 203define ptx_device i32 @test_lanemask_gt() {204; CHECK: mov.u32 %r{{[0-9]+}}, %lanemask_gt;205; CHECK: ret;206 %x = call i32 @llvm.nvvm.read.ptx.sreg.lanemask.gt()207 ret i32 %x208}209 210define ptx_device i32 @test_clock() {211; CHECK: mov.u32 %r{{[0-9]+}}, %clock;212; CHECK: ret;213 %x = call i32 @llvm.nvvm.read.ptx.sreg.clock()214 ret i32 %x215}216 217define ptx_device i64 @test_clock64() {218; CHECK: mov.u64 %rd{{[0-9]+}}, %clock64;219; CHECK: ret;220 %x = call i64 @llvm.nvvm.read.ptx.sreg.clock64()221 ret i64 %x222}223 224define ptx_device i32 @test_pm0() {225; CHECK: mov.u32 %r{{[0-9]+}}, %pm0;226; CHECK: ret;227 %x = call i32 @llvm.nvvm.read.ptx.sreg.pm0()228 ret i32 %x229}230 231define ptx_device i32 @test_pm1() {232; CHECK: mov.u32 %r{{[0-9]+}}, %pm1;233; CHECK: ret;234 %x = call i32 @llvm.nvvm.read.ptx.sreg.pm1()235 ret i32 %x236}237 238define ptx_device i32 @test_pm2() {239; CHECK: mov.u32 %r{{[0-9]+}}, %pm2;240; CHECK: ret;241 %x = call i32 @llvm.nvvm.read.ptx.sreg.pm2()242 ret i32 %x243}244 245define ptx_device i32 @test_pm3() {246; CHECK: mov.u32 %r{{[0-9]+}}, %pm3;247; CHECK: ret;248 %x = call i32 @llvm.nvvm.read.ptx.sreg.pm3()249 ret i32 %x250}251 252define ptx_device void @test_bar_sync() {253; CHECK: bar.sync 0254; CHECK: ret;255 call void @llvm.nvvm.bar.sync(i32 0)256 ret void257}258 259declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()260declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()261declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()262declare i32 @llvm.nvvm.read.ptx.sreg.tid.w()263declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()264declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()265declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()266declare i32 @llvm.nvvm.read.ptx.sreg.ntid.w()267 268declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()269declare i32 @llvm.nvvm.read.ptx.sreg.laneid()270declare i32 @llvm.nvvm.read.ptx.sreg.warpid()271declare i32 @llvm.nvvm.read.ptx.sreg.nwarpid()272 273declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()274declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()275declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()276declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.w()277declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()278declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()279declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()280declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.w()281 282declare i32 @llvm.nvvm.read.ptx.sreg.smid()283declare i32 @llvm.nvvm.read.ptx.sreg.nsmid()284declare i32 @llvm.nvvm.read.ptx.sreg.gridid()285 286declare i32 @llvm.nvvm.read.ptx.sreg.lanemask.eq()287declare i32 @llvm.nvvm.read.ptx.sreg.lanemask.le()288declare i32 @llvm.nvvm.read.ptx.sreg.lanemask.lt()289declare i32 @llvm.nvvm.read.ptx.sreg.lanemask.ge()290declare i32 @llvm.nvvm.read.ptx.sreg.lanemask.gt()291 292declare i32 @llvm.nvvm.read.ptx.sreg.clock()293declare i64 @llvm.nvvm.read.ptx.sreg.clock64()294 295declare i32 @llvm.nvvm.read.ptx.sreg.pm0()296declare i32 @llvm.nvvm.read.ptx.sreg.pm1()297declare i32 @llvm.nvvm.read.ptx.sreg.pm2()298declare i32 @llvm.nvvm.read.ptx.sreg.pm3()299 300declare void @llvm.nvvm.bar.sync(i32 %i)301 302!0 = !{i32 0, i32 19}303