143 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}6 7target triple = "nvptx-unknown-cuda"8 9@tex0 = internal addrspace(1) global i64 0, align 810@surf0 = internal addrspace(1) global i64 0, align 811 12declare i32 @llvm.nvvm.txq.width(i64)13declare i32 @llvm.nvvm.txq.height(i64)14declare i32 @llvm.nvvm.suq.width(i64)15declare i32 @llvm.nvvm.suq.height(i64)16declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))17 18 19define i32 @t0(i64 %texHandle) {20; CHECK-LABEL: t0(21; CHECK: {22; CHECK-NEXT: .reg .b32 %r<2>;23; CHECK-NEXT: .reg .b64 %rd<2>;24; CHECK-EMPTY:25; CHECK-NEXT: // %bb.0:26; CHECK-NEXT: ld.param.b64 %rd1, [t0_param_0];27; CHECK-NEXT: txq.width.b32 %r1, [%rd1];28; CHECK-NEXT: st.param.b32 [func_retval0], %r1;29; CHECK-NEXT: ret;30 %width = tail call i32 @llvm.nvvm.txq.width(i64 %texHandle)31 ret i32 %width32}33 34define i32 @t1() {35; CHECK-LABEL: t1(36; CHECK: {37; CHECK-NEXT: .reg .b32 %r<2>;38; CHECK-EMPTY:39; CHECK-NEXT: // %bb.0:40; CHECK-NEXT: txq.width.b32 %r1, [tex0];41; CHECK-NEXT: st.param.b32 [func_retval0], %r1;42; CHECK-NEXT: ret;43 %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)44 %width = tail call i32 @llvm.nvvm.txq.width(i64 %texHandle)45 ret i32 %width46}47 48 49define i32 @t2(i64 %texHandle) {50; CHECK-LABEL: t2(51; CHECK: {52; CHECK-NEXT: .reg .b32 %r<2>;53; CHECK-NEXT: .reg .b64 %rd<2>;54; CHECK-EMPTY:55; CHECK-NEXT: // %bb.0:56; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_0];57; CHECK-NEXT: txq.height.b32 %r1, [%rd1];58; CHECK-NEXT: st.param.b32 [func_retval0], %r1;59; CHECK-NEXT: ret;60 %height = tail call i32 @llvm.nvvm.txq.height(i64 %texHandle)61 ret i32 %height62}63 64define i32 @t3() {65; CHECK-LABEL: t3(66; CHECK: {67; CHECK-NEXT: .reg .b32 %r<2>;68; CHECK-EMPTY:69; CHECK-NEXT: // %bb.0:70; CHECK-NEXT: txq.height.b32 %r1, [tex0];71; CHECK-NEXT: st.param.b32 [func_retval0], %r1;72; CHECK-NEXT: ret;73 %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)74 %height = tail call i32 @llvm.nvvm.txq.height(i64 %texHandle)75 ret i32 %height76}77 78 79define i32 @s0(i64 %surfHandle) {80; CHECK-LABEL: s0(81; CHECK: {82; CHECK-NEXT: .reg .b32 %r<2>;83; CHECK-NEXT: .reg .b64 %rd<2>;84; CHECK-EMPTY:85; CHECK-NEXT: // %bb.0:86; CHECK-NEXT: ld.param.b64 %rd1, [s0_param_0];87; CHECK-NEXT: suq.width.b32 %r1, [%rd1];88; CHECK-NEXT: st.param.b32 [func_retval0], %r1;89; CHECK-NEXT: ret;90 %width = tail call i32 @llvm.nvvm.suq.width(i64 %surfHandle)91 ret i32 %width92}93 94define i32 @s1() {95; CHECK-LABEL: s1(96; CHECK: {97; CHECK-NEXT: .reg .b32 %r<2>;98; CHECK-EMPTY:99; CHECK-NEXT: // %bb.0:100; CHECK-NEXT: suq.width.b32 %r1, [surf0];101; CHECK-NEXT: st.param.b32 [func_retval0], %r1;102; CHECK-NEXT: ret;103 %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)104 %width = tail call i32 @llvm.nvvm.suq.width(i64 %surfHandle)105 ret i32 %width106}107 108 109define i32 @s2(i64 %surfHandle) {110; CHECK-LABEL: s2(111; CHECK: {112; CHECK-NEXT: .reg .b32 %r<2>;113; CHECK-NEXT: .reg .b64 %rd<2>;114; CHECK-EMPTY:115; CHECK-NEXT: // %bb.0:116; CHECK-NEXT: ld.param.b64 %rd1, [s2_param_0];117; CHECK-NEXT: suq.height.b32 %r1, [%rd1];118; CHECK-NEXT: st.param.b32 [func_retval0], %r1;119; CHECK-NEXT: ret;120 %height = tail call i32 @llvm.nvvm.suq.height(i64 %surfHandle)121 ret i32 %height122}123 124define i32 @s3() {125; CHECK-LABEL: s3(126; CHECK: {127; CHECK-NEXT: .reg .b32 %r<2>;128; CHECK-EMPTY:129; CHECK-NEXT: // %bb.0:130; CHECK-NEXT: suq.height.b32 %r1, [surf0];131; CHECK-NEXT: st.param.b32 [func_retval0], %r1;132; CHECK-NEXT: ret;133 %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)134 %height = tail call i32 @llvm.nvvm.suq.height(i64 %surfHandle)135 ret i32 %height136}137 138 139 140!nvvm.annotations = !{!1, !2}141!1 = !{ptr addrspace(1) @tex0, !"texture", i32 1}142!2 = !{ptr addrspace(1) @surf0, !"surface", i32 1}143