116 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx | FileCheck %s3; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -verify-machineinstrs | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7%struct.float2 = type { float, float }8 9; CHECK-LABEL: .visible .func (.param .b32 func_retval0) callee_md10; CHECK-NEXT: (11; CHECK-NEXT: .param .align 8 .b8 callee_md_param_0[8]12; CHECK-NEXT: )13; CHECK-NEXT: ;14 15; CHECK-LABEL: .visible .func (.param .b32 func_retval0) callee16; CHECK-NEXT: (17; CHECK-NEXT: .param .align 8 .b8 callee_param_0[8]18; CHECK-NEXT: )19; CHECK-NEXT: ;20 21define float @caller_md(float %a, float %b) {22; CHECK-LABEL: caller_md(23; CHECK: {24; CHECK-NEXT: .reg .b32 %r<4>;25; CHECK-EMPTY:26; CHECK-NEXT: // %bb.0:27; CHECK-NEXT: ld.param.b32 %r1, [caller_md_param_0];28; CHECK-NEXT: ld.param.b32 %r2, [caller_md_param_1];29; CHECK-NEXT: { // callseq 0, 030; CHECK-NEXT: .param .align 8 .b8 param0[8];31; CHECK-NEXT: .param .b32 retval0;32; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, %r2};33; CHECK-NEXT: call.uni (retval0), callee_md, (param0);34; CHECK-NEXT: ld.param.b32 %r3, [retval0];35; CHECK-NEXT: } // callseq 036; CHECK-NEXT: st.param.b32 [func_retval0], %r3;37; CHECK-NEXT: ret;38 39 %s1 = insertvalue %struct.float2 poison, float %a, 040 %s2 = insertvalue %struct.float2 %s1, float %b, 141 %r = call float @callee_md(%struct.float2 %s2)42 ret float %r43}44 45define float @callee_md(%struct.float2 alignstack(8) %a) {46; CHECK-LABEL: callee_md(47; CHECK: {48; CHECK-NEXT: .reg .b32 %r<4>;49; CHECK-EMPTY:50; CHECK-NEXT: // %bb.0:51; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [callee_md_param_0];52; CHECK-NEXT: add.rn.f32 %r3, %r1, %r2;53; CHECK-NEXT: st.param.b32 [func_retval0], %r3;54; CHECK-NEXT: ret;55 56 %v0 = extractvalue %struct.float2 %a, 057 %v1 = extractvalue %struct.float2 %a, 158 %2 = fadd float %v0, %v159 ret float %260}61 62define float @caller(float %a, float %b) {63; CHECK-LABEL: caller(64; CHECK: {65; CHECK-NEXT: .reg .b32 %r<4>;66; CHECK-EMPTY:67; CHECK-NEXT: // %bb.0:68; CHECK-NEXT: ld.param.b32 %r1, [caller_param_0];69; CHECK-NEXT: ld.param.b32 %r2, [caller_param_1];70; CHECK-NEXT: { // callseq 1, 071; CHECK-NEXT: .param .align 8 .b8 param0[8];72; CHECK-NEXT: .param .b32 retval0;73; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, %r2};74; CHECK-NEXT: call.uni (retval0), callee, (param0);75; CHECK-NEXT: ld.param.b32 %r3, [retval0];76; CHECK-NEXT: } // callseq 177; CHECK-NEXT: st.param.b32 [func_retval0], %r3;78; CHECK-NEXT: ret;79 80 %s1 = insertvalue %struct.float2 poison, float %a, 081 %s2 = insertvalue %struct.float2 %s1, float %b, 182 %r = call float @callee(%struct.float2 %s2)83 ret float %r84}85 86define float @callee(%struct.float2 alignstack(8) %a ) {87; CHECK-LABEL: callee(88; CHECK: {89; CHECK-NEXT: .reg .b32 %r<4>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [callee_param_0];93; CHECK-NEXT: add.rn.f32 %r3, %r1, %r2;94; CHECK-NEXT: st.param.b32 [func_retval0], %r3;95; CHECK-NEXT: ret;96 97 %v0 = extractvalue %struct.float2 %a, 098 %v1 = extractvalue %struct.float2 %a, 199 %2 = fadd float %v0, %v1100 ret float %2101}102 103define alignstack(8) %struct.float2 @aligned_return(%struct.float2 %a ) {104; CHECK-LABEL: aligned_return(105; CHECK: {106; CHECK-NEXT: .reg .b32 %r<3>;107; CHECK-EMPTY:108; CHECK-NEXT: // %bb.0:109; CHECK-NEXT: ld.param.b32 %r1, [aligned_return_param_0];110; CHECK-NEXT: ld.param.b32 %r2, [aligned_return_param_0+4];111; CHECK-NEXT: st.param.b32 [func_retval0+4], %r2;112; CHECK-NEXT: st.param.b32 [func_retval0], %r1;113; CHECK-NEXT: ret;114 ret %struct.float2 %a115}116