91 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 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<6>;17; CHECK-NEXT: .reg .b64 %rd<4>;18; CHECK-EMPTY:19; CHECK-NEXT: // %bb.0:20; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];21; CHECK-NEXT: ld.param.b64 %rd2, [foo_param_1];22; CHECK-NEXT: cvta.to.global.u64 %rd3, %rd2;23; CHECK-NEXT: ld.param.b32 %r1, [foo_param_2];24; CHECK-NEXT: tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [%rd1, {%r1}];25; CHECK-NEXT: st.global.b32 [%rd3], %r2;26; CHECK-NEXT: ret;27 %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx)28 %ret = extractvalue { float, float, float, float } %val, 029 store float %ret, ptr %red30 ret void31}32 33 34@tex0 = internal addrspace(1) global i64 0, align 835 36define ptx_kernel void @bar(ptr %red, i32 %idx) {37; CHECK-LABEL: bar(38; CHECK: {39; CHECK-NEXT: .reg .b32 %r<6>;40; CHECK-NEXT: .reg .b64 %rd<3>;41; CHECK-EMPTY:42; CHECK-NEXT: // %bb.0:43; CHECK-NEXT: ld.param.b64 %rd1, [bar_param_0];44; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;45; CHECK-NEXT: ld.param.b32 %r1, [bar_param_1];46; CHECK-NEXT: tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [tex0, {%r1}];47; CHECK-NEXT: st.global.b32 [%rd2], %r2;48; CHECK-NEXT: ret;49 %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)50 %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)51 %ret = extractvalue { float, float, float, float } %val, 052 store float %ret, ptr %red53 ret void54}55 56declare float @texfunc(i64)57 58define ptx_kernel void @baz(ptr %red, i32 %idx) {59; CHECK-LABEL: baz(60; CHECK: {61; CHECK-NEXT: .reg .b32 %r<8>;62; CHECK-NEXT: .reg .b64 %rd<4>;63; CHECK-EMPTY:64; CHECK-NEXT: // %bb.0:65; CHECK-NEXT: ld.param.b64 %rd1, [baz_param_0];66; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;67; CHECK-NEXT: ld.param.b32 %r1, [baz_param_1];68; CHECK-NEXT: mov.u64 %rd3, tex0;69; CHECK-NEXT: tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [tex0, {%r1}];70; CHECK-NEXT: { // callseq 0, 071; CHECK-NEXT: .param .b64 param0;72; CHECK-NEXT: .param .b32 retval0;73; CHECK-NEXT: st.param.b64 [param0], %rd3;74; CHECK-NEXT: call.uni (retval0), texfunc, (param0);75; CHECK-NEXT: ld.param.b32 %r6, [retval0];76; CHECK-NEXT: } // callseq 077; CHECK-NEXT: add.rn.f32 %r7, %r2, %r6;78; CHECK-NEXT: st.global.b32 [%rd2], %r7;79; CHECK-NEXT: ret;80 %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)81 %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)82 %ret = extractvalue { float, float, float, float } %val, 083 %texcall = tail call float @texfunc(i64 %texHandle)84 %ret2 = fadd float %ret, %texcall85 store float %ret2, ptr %red86 ret void87}88 89!nvvm.annotations = !{!1}90!1 = !{ptr addrspace(1) @tex0, !"texture", i32 1}91