118 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_35 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}4 5declare <2 x float> @barv(<2 x float> %input)6declare <3 x float> @barv3(<3 x float> %input)7declare [2 x float] @bara([2 x float] %input)8declare {float, float} @bars({float, float} %input)9 10define void @test_v2f32(<2 x float> %input, ptr %output) {11; CHECK-LABEL: test_v2f32(12; CHECK: {13; CHECK-NEXT: .reg .b32 %r<5>;14; CHECK-NEXT: .reg .b64 %rd<2>;15; CHECK-EMPTY:16; CHECK-NEXT: // %bb.0:17; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_v2f32_param_0];18; CHECK-NEXT: { // callseq 0, 019; CHECK-NEXT: .param .align 8 .b8 param0[8];20; CHECK-NEXT: .param .align 8 .b8 retval0[8];21; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, %r2};22; CHECK-NEXT: call.uni (retval0), barv, (param0);23; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [retval0];24; CHECK-NEXT: } // callseq 025; CHECK-NEXT: ld.param.b64 %rd1, [test_v2f32_param_1];26; CHECK-NEXT: st.v2.b32 [%rd1], {%r3, %r4};27; CHECK-NEXT: ret;28 %call = tail call <2 x float> @barv(<2 x float> %input)29 store <2 x float> %call, ptr %output, align 830 ret void31}32 33define void @test_v3f32(<3 x float> %input, ptr %output) {34; CHECK-LABEL: test_v3f32(35; CHECK: {36; CHECK-NEXT: .reg .b32 %r<7>;37; CHECK-NEXT: .reg .b64 %rd<6>;38; CHECK-EMPTY:39; CHECK-NEXT: // %bb.0:40; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_v3f32_param_0];41; CHECK-NEXT: ld.param.b32 %r3, [test_v3f32_param_0+8];42; CHECK-NEXT: { // callseq 1, 043; CHECK-NEXT: .param .align 16 .b8 param0[16];44; CHECK-NEXT: .param .align 16 .b8 retval0[16];45; CHECK-NEXT: st.param.b32 [param0+8], %r3;46; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, %r2};47; CHECK-NEXT: call.uni (retval0), barv3, (param0);48; CHECK-NEXT: ld.param.b32 %r4, [retval0+8];49; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];50; CHECK-NEXT: } // callseq 151; CHECK-NEXT: cvt.u64.u32 %rd1, %r5;52; CHECK-NEXT: cvt.u64.u32 %rd2, %r6;53; CHECK-NEXT: shl.b64 %rd3, %rd2, 32;54; CHECK-NEXT: or.b64 %rd4, %rd1, %rd3;55; CHECK-NEXT: ld.param.b64 %rd5, [test_v3f32_param_1];56; CHECK-NEXT: st.b32 [%rd5+8], %r4;57; CHECK-NEXT: st.b64 [%rd5], %rd4;58; CHECK-NEXT: ret;59 %call = tail call <3 x float> @barv3(<3 x float> %input)60; Make sure we don't load more values than than we need to.61 store <3 x float> %call, ptr %output, align 862 ret void63}64 65define void @test_a2f32([2 x float] %input, ptr %output) {66; CHECK-LABEL: test_a2f32(67; CHECK: {68; CHECK-NEXT: .reg .b32 %r<5>;69; CHECK-NEXT: .reg .b64 %rd<2>;70; CHECK-EMPTY:71; CHECK-NEXT: // %bb.0:72; CHECK-NEXT: ld.param.b32 %r1, [test_a2f32_param_0];73; CHECK-NEXT: ld.param.b32 %r2, [test_a2f32_param_0+4];74; CHECK-NEXT: { // callseq 2, 075; CHECK-NEXT: .param .align 4 .b8 param0[8];76; CHECK-NEXT: .param .align 4 .b8 retval0[8];77; CHECK-NEXT: st.param.b32 [param0+4], %r2;78; CHECK-NEXT: st.param.b32 [param0], %r1;79; CHECK-NEXT: call.uni (retval0), bara, (param0);80; CHECK-NEXT: ld.param.b32 %r3, [retval0+4];81; CHECK-NEXT: ld.param.b32 %r4, [retval0];82; CHECK-NEXT: } // callseq 283; CHECK-NEXT: ld.param.b64 %rd1, [test_a2f32_param_1];84; CHECK-NEXT: st.b32 [%rd1+4], %r3;85; CHECK-NEXT: st.b32 [%rd1], %r4;86; CHECK-NEXT: ret;87 %call = tail call [2 x float] @bara([2 x float] %input)88 store [2 x float] %call, ptr %output, align 489 ret void90}91 92define void @test_s2f32({float, float} %input, ptr %output) {93; CHECK-LABEL: test_s2f32(94; CHECK: {95; CHECK-NEXT: .reg .b32 %r<5>;96; CHECK-NEXT: .reg .b64 %rd<2>;97; CHECK-EMPTY:98; CHECK-NEXT: // %bb.0:99; CHECK-NEXT: ld.param.b32 %r1, [test_s2f32_param_0];100; CHECK-NEXT: ld.param.b32 %r2, [test_s2f32_param_0+4];101; CHECK-NEXT: { // callseq 3, 0102; CHECK-NEXT: .param .align 4 .b8 param0[8];103; CHECK-NEXT: .param .align 4 .b8 retval0[8];104; CHECK-NEXT: st.param.b32 [param0+4], %r2;105; CHECK-NEXT: st.param.b32 [param0], %r1;106; CHECK-NEXT: call.uni (retval0), bars, (param0);107; CHECK-NEXT: ld.param.b32 %r3, [retval0+4];108; CHECK-NEXT: ld.param.b32 %r4, [retval0];109; CHECK-NEXT: } // callseq 3110; CHECK-NEXT: ld.param.b64 %rd1, [test_s2f32_param_1];111; CHECK-NEXT: st.b32 [%rd1+4], %r3;112; CHECK-NEXT: st.b32 [%rd1], %r4;113; CHECK-NEXT: ret;114 %call = tail call {float, float} @bars({float, float} %input)115 store {float, float} %call, ptr %output, align 4116 ret void117}118