1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s 3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s 4; 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 7 8target triple = "nvptx-unknown-cuda" 9 10declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64, i32) 11declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1)) 12 13define ptx_kernel void @foo(i64 %img, ptr %red, i32 %idx) { 14; CHECK-LABEL: foo( 15; CHECK: { 16; CHECK-NEXT: .reg .b32 %r<2>; 17; CHECK-NEXT: .reg .f32 %f<5>; 18; CHECK-NEXT: .reg .b64 %rd<4>; 19; CHECK-EMPTY: 20; CHECK-NEXT: // %bb.0: 21; CHECK-NEXT: ld.param.u64 %rd1, [foo_param_0]; 22; CHECK-NEXT: ld.param.u64 %rd2, [foo_param_1]; 23; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2; 24; CHECK-NEXT: ld.param.u32 %r1, [foo_param_2]; 25; CHECK-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [%rd1, {%r1}]; 26; CHECK-NEXT: st.global.f32 [%rd3], %f1; 27; CHECK-NEXT: ret; 28 %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx) 29 %ret = extractvalue { float, float, float, float } %val, 0 30 store float %ret, ptr %red 31 ret void 32} 33 34 35@tex0 = internal addrspace(1) global i64 0, align 8 36 37define ptx_kernel void @bar(ptr %red, i32 %idx) { 38; CHECK-LABEL: bar( 39; CHECK: { 40; CHECK-NEXT: .reg .b32 %r<2>; 41; CHECK-NEXT: .reg .f32 %f<5>; 42; CHECK-NEXT: .reg .b64 %rd<4>; 43; CHECK-EMPTY: 44; CHECK-NEXT: // %bb.0: 45; CHECK-NEXT: ld.param.u64 %rd1, [bar_param_0]; 46; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1; 47; CHECK-NEXT: ld.param.u32 %r1, [bar_param_1]; 48; CHECK-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}]; 49; CHECK-NEXT: st.global.f32 [%rd2], %f1; 50; CHECK-NEXT: ret; 51 %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0) 52 %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx) 53 %ret = extractvalue { float, float, float, float } %val, 0 54 store float %ret, ptr %red 55 ret void 56} 57 58declare float @texfunc(i64) 59 60define ptx_kernel void @baz(ptr %red, i32 %idx) { 61; CHECK-LABEL: baz( 62; CHECK: { 63; CHECK-NEXT: .reg .b32 %r<2>; 64; CHECK-NEXT: .reg .f32 %f<8>; 65; CHECK-NEXT: .reg .b64 %rd<4>; 66; CHECK-EMPTY: 67; CHECK-NEXT: // %bb.0: 68; CHECK-NEXT: ld.param.u64 %rd1, [baz_param_0]; 69; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1; 70; CHECK-NEXT: ld.param.u32 %r1, [baz_param_1]; 71; CHECK-NEXT: mov.u64 %rd3, tex0; 72; CHECK-NEXT: tex.1d.v4.f32.s32 {%f1, %f2, %f3, %f4}, [tex0, {%r1}]; 73; CHECK-NEXT: { // callseq 0, 0 74; CHECK-NEXT: .param .b64 param0; 75; CHECK-NEXT: st.param.b64 [param0], %rd3; 76; CHECK-NEXT: .param .b32 retval0; 77; CHECK-NEXT: call.uni (retval0), 78; CHECK-NEXT: texfunc, 79; CHECK-NEXT: ( 80; CHECK-NEXT: param0 81; CHECK-NEXT: ); 82; CHECK-NEXT: ld.param.f32 %f5, [retval0]; 83; CHECK-NEXT: } // callseq 0 84; CHECK-NEXT: add.rn.f32 %f7, %f1, %f5; 85; CHECK-NEXT: st.global.f32 [%rd2], %f7; 86; CHECK-NEXT: ret; 87 %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0) 88 %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx) 89 %ret = extractvalue { float, float, float, float } %val, 0 90 %texcall = tail call float @texfunc(i64 %texHandle) 91 %ret2 = fadd float %ret, %texcall 92 store float %ret2, ptr %red 93 ret void 94} 95 96!nvvm.annotations = !{!1} 97!1 = !{ptr addrspace(1) @tex0, !"texture", i32 1} 98