399 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; Verifies correctness of load/store of parameters and return values.3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck -allow-deprecated-dag-overlap %s4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | %ptxas-verify %}5 6%s_i8i16p = type { <{ i16, i8, i16 }>, i64 }7%s_i8i32p = type { <{ i32, i8, i32 }>, i64 }8%s_i8i64p = type { <{ i64, i8, i64 }>, i64 }9%s_i8f16p = type { <{ half, i8, half }>, i64 }10%s_i8f16x2p = type { <{ <2 x half>, i8, <2 x half> }>, i64 }11%s_i8f32p = type { <{ float, i8, float }>, i64 }12%s_i8f64p = type { <{ double, i8, double }>, i64 }13 14; -- All loads/stores from parameters aligned by one must be done one15; byte at a time.16; -- Notes:17; -- There are two fields of interest in the packed part of the struct, one18; with a proper offset and one without. The former should be loaded or19; stored as a whole, and the latter by bytes.20; -- Only loading and storing the said fields are checked in the following21; series of tests so that they are more concise.22 23 24define %s_i8i16p @test_s_i8i16p(%s_i8i16p %a) {25; CHECK-LABEL: test_s_i8i16p(26; CHECK: {27; CHECK-NEXT: .reg .b16 %rs<9>;28; CHECK-NEXT: .reg .b32 %r<2>;29; CHECK-NEXT: .reg .b64 %rd<3>;30; CHECK-EMPTY:31; CHECK-NEXT: // %bb.0:32; CHECK-NEXT: ld.param.b32 %r1, [test_s_i8i16p_param_0];33; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8i16p_param_0+8];34; CHECK-NEXT: ld.param.b8 %rs1, [test_s_i8i16p_param_0+4];35; CHECK-NEXT: { // callseq 0, 036; CHECK-NEXT: .param .align 8 .b8 param0[16];37; CHECK-NEXT: .param .align 8 .b8 retval0[16];38; CHECK-NEXT: st.param.b8 [param0+4], %rs1;39; CHECK-NEXT: st.param.b64 [param0+8], %rd1;40; CHECK-NEXT: st.param.b32 [param0], %r1;41; CHECK-NEXT: call.uni (retval0), test_s_i8i16p, (param0);42; CHECK-NEXT: ld.param.b64 %rd2, [retval0+8];43; CHECK-NEXT: ld.param.b8 %rs2, [retval0+2];44; CHECK-NEXT: ld.param.b16 %rs3, [retval0];45; CHECK-NEXT: ld.param.b8 %rs4, [retval0+4];46; CHECK-NEXT: ld.param.b8 %rs5, [retval0+3];47; CHECK-NEXT: } // callseq 048; CHECK-NEXT: shl.b16 %rs6, %rs4, 8;49; CHECK-NEXT: or.b16 %rs7, %rs6, %rs5;50; CHECK-NEXT: st.param.b8 [func_retval0+3], %rs5;51; CHECK-NEXT: st.param.b64 [func_retval0+8], %rd2;52; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs2;53; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;54; CHECK-NEXT: shr.u16 %rs8, %rs7, 8;55; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs8;56; CHECK-NEXT: ret;57 %r = tail call %s_i8i16p @test_s_i8i16p(%s_i8i16p %a)58 ret %s_i8i16p %r59}60 61 62define %s_i8i32p @test_s_i8i32p(%s_i8i32p %a) {63; CHECK-LABEL: test_s_i8i32p(64; CHECK: {65; CHECK-NEXT: .reg .b16 %rs<3>;66; CHECK-NEXT: .reg .b32 %r<19>;67; CHECK-NEXT: .reg .b64 %rd<3>;68; CHECK-EMPTY:69; CHECK-NEXT: // %bb.0:70; CHECK-NEXT: ld.param.b32 %r1, [test_s_i8i32p_param_0];71; CHECK-NEXT: ld.param.b16 %rs1, [test_s_i8i32p_param_0+4];72; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8i32p_param_0+16];73; CHECK-NEXT: ld.param.b8 %r2, [test_s_i8i32p_param_0+6];74; CHECK-NEXT: ld.param.b8 %r3, [test_s_i8i32p_param_0+7];75; CHECK-NEXT: ld.param.b8 %r4, [test_s_i8i32p_param_0+8];76; CHECK-NEXT: { // callseq 1, 077; CHECK-NEXT: .param .align 8 .b8 param0[24];78; CHECK-NEXT: .param .align 8 .b8 retval0[24];79; CHECK-NEXT: st.param.b8 [param0+8], %r4;80; CHECK-NEXT: st.param.b8 [param0+7], %r3;81; CHECK-NEXT: st.param.b8 [param0+6], %r2;82; CHECK-NEXT: st.param.b64 [param0+16], %rd1;83; CHECK-NEXT: st.param.b16 [param0+4], %rs1;84; CHECK-NEXT: st.param.b32 [param0], %r1;85; CHECK-NEXT: call.uni (retval0), test_s_i8i32p, (param0);86; CHECK-NEXT: ld.param.b64 %rd2, [retval0+16];87; CHECK-NEXT: ld.param.b8 %rs2, [retval0+4];88; CHECK-NEXT: ld.param.b32 %r5, [retval0];89; CHECK-NEXT: ld.param.b8 %r6, [retval0+8];90; CHECK-NEXT: ld.param.b8 %r7, [retval0+7];91; CHECK-NEXT: ld.param.b8 %r8, [retval0+6];92; CHECK-NEXT: ld.param.b8 %r9, [retval0+5];93; CHECK-NEXT: } // callseq 194; CHECK-NEXT: shl.b32 %r10, %r8, 8;95; CHECK-NEXT: or.b32 %r11, %r10, %r9;96; CHECK-NEXT: shl.b32 %r12, %r7, 16;97; CHECK-NEXT: shl.b32 %r13, %r6, 24;98; CHECK-NEXT: or.b32 %r14, %r13, %r12;99; CHECK-NEXT: or.b32 %r15, %r14, %r11;100; CHECK-NEXT: st.param.b8 [func_retval0+5], %r9;101; CHECK-NEXT: st.param.b64 [func_retval0+16], %rd2;102; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs2;103; CHECK-NEXT: st.param.b32 [func_retval0], %r5;104; CHECK-NEXT: shr.u32 %r16, %r15, 24;105; CHECK-NEXT: st.param.b8 [func_retval0+8], %r16;106; CHECK-NEXT: shr.u32 %r17, %r15, 16;107; CHECK-NEXT: st.param.b8 [func_retval0+7], %r17;108; CHECK-NEXT: shr.u32 %r18, %r15, 8;109; CHECK-NEXT: st.param.b8 [func_retval0+6], %r18;110; CHECK-NEXT: ret;111 %r = tail call %s_i8i32p @test_s_i8i32p(%s_i8i32p %a)112 ret %s_i8i32p %r113}114 115 116define %s_i8i64p @test_s_i8i64p(%s_i8i64p %a) {117; CHECK-LABEL: test_s_i8i64p(118; CHECK: {119; CHECK-NEXT: .reg .b16 %rs<2>;120; CHECK-NEXT: .reg .b64 %rd<36>;121; CHECK-EMPTY:122; CHECK-NEXT: // %bb.0:123; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8i64p_param_0];124; CHECK-NEXT: ld.param.b64 %rd2, [test_s_i8i64p_param_0+8];125; CHECK-NEXT: ld.param.b64 %rd3, [test_s_i8i64p_param_0+24];126; CHECK-NEXT: ld.param.b8 %rd4, [test_s_i8i64p_param_0+16];127; CHECK-NEXT: { // callseq 2, 0128; CHECK-NEXT: .param .align 8 .b8 param0[32];129; CHECK-NEXT: .param .align 8 .b8 retval0[32];130; CHECK-NEXT: st.param.b8 [param0+16], %rd4;131; CHECK-NEXT: st.param.b64 [param0+24], %rd3;132; CHECK-NEXT: st.param.b64 [param0+8], %rd2;133; CHECK-NEXT: st.param.b64 [param0], %rd1;134; CHECK-NEXT: call.uni (retval0), test_s_i8i64p, (param0);135; CHECK-NEXT: ld.param.b64 %rd5, [retval0+24];136; CHECK-NEXT: ld.param.b8 %rs1, [retval0+8];137; CHECK-NEXT: ld.param.b64 %rd6, [retval0];138; CHECK-NEXT: ld.param.b8 %rd7, [retval0+16];139; CHECK-NEXT: ld.param.b8 %rd8, [retval0+15];140; CHECK-NEXT: ld.param.b8 %rd9, [retval0+14];141; CHECK-NEXT: ld.param.b8 %rd10, [retval0+13];142; CHECK-NEXT: ld.param.b8 %rd11, [retval0+12];143; CHECK-NEXT: ld.param.b8 %rd12, [retval0+11];144; CHECK-NEXT: ld.param.b8 %rd13, [retval0+10];145; CHECK-NEXT: ld.param.b8 %rd14, [retval0+9];146; CHECK-NEXT: } // callseq 2147; CHECK-NEXT: shl.b64 %rd15, %rd13, 8;148; CHECK-NEXT: or.b64 %rd16, %rd15, %rd14;149; CHECK-NEXT: shl.b64 %rd17, %rd12, 16;150; CHECK-NEXT: shl.b64 %rd18, %rd11, 24;151; CHECK-NEXT: or.b64 %rd19, %rd18, %rd17;152; CHECK-NEXT: or.b64 %rd20, %rd19, %rd16;153; CHECK-NEXT: shl.b64 %rd21, %rd9, 8;154; CHECK-NEXT: or.b64 %rd22, %rd21, %rd10;155; CHECK-NEXT: shl.b64 %rd23, %rd8, 16;156; CHECK-NEXT: shl.b64 %rd24, %rd7, 24;157; CHECK-NEXT: or.b64 %rd25, %rd24, %rd23;158; CHECK-NEXT: or.b64 %rd26, %rd25, %rd22;159; CHECK-NEXT: shl.b64 %rd27, %rd26, 32;160; CHECK-NEXT: or.b64 %rd28, %rd27, %rd20;161; CHECK-NEXT: st.param.b8 [func_retval0+9], %rd14;162; CHECK-NEXT: st.param.b64 [func_retval0+24], %rd5;163; CHECK-NEXT: st.param.b8 [func_retval0+8], %rs1;164; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;165; CHECK-NEXT: shr.u64 %rd29, %rd28, 56;166; CHECK-NEXT: st.param.b8 [func_retval0+16], %rd29;167; CHECK-NEXT: shr.u64 %rd30, %rd28, 48;168; CHECK-NEXT: st.param.b8 [func_retval0+15], %rd30;169; CHECK-NEXT: shr.u64 %rd31, %rd28, 40;170; CHECK-NEXT: st.param.b8 [func_retval0+14], %rd31;171; CHECK-NEXT: shr.u64 %rd32, %rd28, 32;172; CHECK-NEXT: st.param.b8 [func_retval0+13], %rd32;173; CHECK-NEXT: shr.u64 %rd33, %rd28, 24;174; CHECK-NEXT: st.param.b8 [func_retval0+12], %rd33;175; CHECK-NEXT: shr.u64 %rd34, %rd28, 16;176; CHECK-NEXT: st.param.b8 [func_retval0+11], %rd34;177; CHECK-NEXT: shr.u64 %rd35, %rd28, 8;178; CHECK-NEXT: st.param.b8 [func_retval0+10], %rd35;179; CHECK-NEXT: ret;180 %r = tail call %s_i8i64p @test_s_i8i64p(%s_i8i64p %a)181 ret %s_i8i64p %r182}183 184 185define %s_i8f16p @test_s_i8f16p(%s_i8f16p %a) {186; CHECK-LABEL: test_s_i8f16p(187; CHECK: {188; CHECK-NEXT: .reg .b16 %rs<11>;189; CHECK-NEXT: .reg .b64 %rd<3>;190; CHECK-EMPTY:191; CHECK-NEXT: // %bb.0:192; CHECK-NEXT: ld.param.b16 %rs1, [test_s_i8f16p_param_0];193; CHECK-NEXT: ld.param.b16 %rs2, [test_s_i8f16p_param_0+2];194; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8f16p_param_0+8];195; CHECK-NEXT: ld.param.b8 %rs3, [test_s_i8f16p_param_0+4];196; CHECK-NEXT: { // callseq 3, 0197; CHECK-NEXT: .param .align 8 .b8 param0[16];198; CHECK-NEXT: .param .align 8 .b8 retval0[16];199; CHECK-NEXT: st.param.b8 [param0+4], %rs3;200; CHECK-NEXT: st.param.b64 [param0+8], %rd1;201; CHECK-NEXT: st.param.b16 [param0+2], %rs2;202; CHECK-NEXT: st.param.b16 [param0], %rs1;203; CHECK-NEXT: call.uni (retval0), test_s_i8f16p, (param0);204; CHECK-NEXT: ld.param.b64 %rd2, [retval0+8];205; CHECK-NEXT: ld.param.b8 %rs4, [retval0+2];206; CHECK-NEXT: ld.param.b16 %rs5, [retval0];207; CHECK-NEXT: ld.param.b8 %rs6, [retval0+4];208; CHECK-NEXT: ld.param.b8 %rs7, [retval0+3];209; CHECK-NEXT: } // callseq 3210; CHECK-NEXT: shl.b16 %rs8, %rs6, 8;211; CHECK-NEXT: or.b16 %rs9, %rs8, %rs7;212; CHECK-NEXT: st.param.b8 [func_retval0+3], %rs7;213; CHECK-NEXT: st.param.b64 [func_retval0+8], %rd2;214; CHECK-NEXT: st.param.b8 [func_retval0+2], %rs4;215; CHECK-NEXT: st.param.b16 [func_retval0], %rs5;216; CHECK-NEXT: shr.u16 %rs10, %rs9, 8;217; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs10;218; CHECK-NEXT: ret;219 %r = tail call %s_i8f16p @test_s_i8f16p(%s_i8f16p %a)220 ret %s_i8f16p %r221}222 223 224define %s_i8f16x2p @test_s_i8f16x2p(%s_i8f16x2p %a) {225; CHECK-LABEL: test_s_i8f16x2p(226; CHECK: {227; CHECK-NEXT: .reg .b16 %rs<3>;228; CHECK-NEXT: .reg .b32 %r<19>;229; CHECK-NEXT: .reg .b64 %rd<3>;230; CHECK-EMPTY:231; CHECK-NEXT: // %bb.0:232; CHECK-NEXT: ld.param.b32 %r1, [test_s_i8f16x2p_param_0];233; CHECK-NEXT: ld.param.b16 %rs1, [test_s_i8f16x2p_param_0+4];234; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8f16x2p_param_0+16];235; CHECK-NEXT: ld.param.b8 %r2, [test_s_i8f16x2p_param_0+6];236; CHECK-NEXT: ld.param.b8 %r3, [test_s_i8f16x2p_param_0+7];237; CHECK-NEXT: ld.param.b8 %r4, [test_s_i8f16x2p_param_0+8];238; CHECK-NEXT: { // callseq 4, 0239; CHECK-NEXT: .param .align 8 .b8 param0[24];240; CHECK-NEXT: .param .align 8 .b8 retval0[24];241; CHECK-NEXT: st.param.b8 [param0+8], %r4;242; CHECK-NEXT: st.param.b8 [param0+7], %r3;243; CHECK-NEXT: st.param.b8 [param0+6], %r2;244; CHECK-NEXT: st.param.b64 [param0+16], %rd1;245; CHECK-NEXT: st.param.b16 [param0+4], %rs1;246; CHECK-NEXT: st.param.b32 [param0], %r1;247; CHECK-NEXT: call.uni (retval0), test_s_i8f16x2p, (param0);248; CHECK-NEXT: ld.param.b64 %rd2, [retval0+16];249; CHECK-NEXT: ld.param.b8 %rs2, [retval0+4];250; CHECK-NEXT: ld.param.b32 %r5, [retval0];251; CHECK-NEXT: ld.param.b8 %r6, [retval0+8];252; CHECK-NEXT: ld.param.b8 %r7, [retval0+7];253; CHECK-NEXT: ld.param.b8 %r8, [retval0+6];254; CHECK-NEXT: ld.param.b8 %r9, [retval0+5];255; CHECK-NEXT: } // callseq 4256; CHECK-NEXT: shl.b32 %r10, %r8, 8;257; CHECK-NEXT: or.b32 %r11, %r10, %r9;258; CHECK-NEXT: shl.b32 %r12, %r7, 16;259; CHECK-NEXT: shl.b32 %r13, %r6, 24;260; CHECK-NEXT: or.b32 %r14, %r13, %r12;261; CHECK-NEXT: or.b32 %r15, %r14, %r11;262; CHECK-NEXT: st.param.b8 [func_retval0+5], %r9;263; CHECK-NEXT: st.param.b64 [func_retval0+16], %rd2;264; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs2;265; CHECK-NEXT: st.param.b32 [func_retval0], %r5;266; CHECK-NEXT: shr.u32 %r16, %r15, 24;267; CHECK-NEXT: st.param.b8 [func_retval0+8], %r16;268; CHECK-NEXT: shr.u32 %r17, %r15, 16;269; CHECK-NEXT: st.param.b8 [func_retval0+7], %r17;270; CHECK-NEXT: shr.u32 %r18, %r15, 8;271; CHECK-NEXT: st.param.b8 [func_retval0+6], %r18;272; CHECK-NEXT: ret;273 %r = tail call %s_i8f16x2p @test_s_i8f16x2p(%s_i8f16x2p %a)274 ret %s_i8f16x2p %r275}276 277 278define %s_i8f32p @test_s_i8f32p(%s_i8f32p %a) {279; CHECK-LABEL: test_s_i8f32p(280; CHECK: {281; CHECK-NEXT: .reg .b16 %rs<3>;282; CHECK-NEXT: .reg .b32 %r<19>;283; CHECK-NEXT: .reg .b64 %rd<3>;284; CHECK-EMPTY:285; CHECK-NEXT: // %bb.0:286; CHECK-NEXT: ld.param.b32 %r1, [test_s_i8f32p_param_0];287; CHECK-NEXT: ld.param.b16 %rs1, [test_s_i8f32p_param_0+4];288; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8f32p_param_0+16];289; CHECK-NEXT: ld.param.b8 %r2, [test_s_i8f32p_param_0+6];290; CHECK-NEXT: ld.param.b8 %r3, [test_s_i8f32p_param_0+7];291; CHECK-NEXT: ld.param.b8 %r4, [test_s_i8f32p_param_0+8];292; CHECK-NEXT: { // callseq 5, 0293; CHECK-NEXT: .param .align 8 .b8 param0[24];294; CHECK-NEXT: .param .align 8 .b8 retval0[24];295; CHECK-NEXT: st.param.b8 [param0+8], %r4;296; CHECK-NEXT: st.param.b8 [param0+7], %r3;297; CHECK-NEXT: st.param.b8 [param0+6], %r2;298; CHECK-NEXT: st.param.b64 [param0+16], %rd1;299; CHECK-NEXT: st.param.b16 [param0+4], %rs1;300; CHECK-NEXT: st.param.b32 [param0], %r1;301; CHECK-NEXT: call.uni (retval0), test_s_i8f32p, (param0);302; CHECK-NEXT: ld.param.b64 %rd2, [retval0+16];303; CHECK-NEXT: ld.param.b8 %rs2, [retval0+4];304; CHECK-NEXT: ld.param.b32 %r5, [retval0];305; CHECK-NEXT: ld.param.b8 %r6, [retval0+8];306; CHECK-NEXT: ld.param.b8 %r7, [retval0+7];307; CHECK-NEXT: ld.param.b8 %r8, [retval0+6];308; CHECK-NEXT: ld.param.b8 %r9, [retval0+5];309; CHECK-NEXT: } // callseq 5310; CHECK-NEXT: shl.b32 %r10, %r8, 8;311; CHECK-NEXT: or.b32 %r11, %r10, %r9;312; CHECK-NEXT: shl.b32 %r12, %r7, 16;313; CHECK-NEXT: shl.b32 %r13, %r6, 24;314; CHECK-NEXT: or.b32 %r14, %r13, %r12;315; CHECK-NEXT: or.b32 %r15, %r14, %r11;316; CHECK-NEXT: st.param.b8 [func_retval0+5], %r9;317; CHECK-NEXT: st.param.b64 [func_retval0+16], %rd2;318; CHECK-NEXT: st.param.b8 [func_retval0+4], %rs2;319; CHECK-NEXT: st.param.b32 [func_retval0], %r5;320; CHECK-NEXT: shr.u32 %r16, %r15, 24;321; CHECK-NEXT: st.param.b8 [func_retval0+8], %r16;322; CHECK-NEXT: shr.u32 %r17, %r15, 16;323; CHECK-NEXT: st.param.b8 [func_retval0+7], %r17;324; CHECK-NEXT: shr.u32 %r18, %r15, 8;325; CHECK-NEXT: st.param.b8 [func_retval0+6], %r18;326; CHECK-NEXT: ret;327 %r = tail call %s_i8f32p @test_s_i8f32p(%s_i8f32p %a)328 ret %s_i8f32p %r329}330 331 332define %s_i8f64p @test_s_i8f64p(%s_i8f64p %a) {333; CHECK-LABEL: test_s_i8f64p(334; CHECK: {335; CHECK-NEXT: .reg .b16 %rs<2>;336; CHECK-NEXT: .reg .b64 %rd<36>;337; CHECK-EMPTY:338; CHECK-NEXT: // %bb.0:339; CHECK-NEXT: ld.param.b64 %rd1, [test_s_i8f64p_param_0];340; CHECK-NEXT: ld.param.b64 %rd2, [test_s_i8f64p_param_0+8];341; CHECK-NEXT: ld.param.b64 %rd3, [test_s_i8f64p_param_0+24];342; CHECK-NEXT: ld.param.b8 %rd4, [test_s_i8f64p_param_0+16];343; CHECK-NEXT: { // callseq 6, 0344; CHECK-NEXT: .param .align 8 .b8 param0[32];345; CHECK-NEXT: .param .align 8 .b8 retval0[32];346; CHECK-NEXT: st.param.b8 [param0+16], %rd4;347; CHECK-NEXT: st.param.b64 [param0+24], %rd3;348; CHECK-NEXT: st.param.b64 [param0+8], %rd2;349; CHECK-NEXT: st.param.b64 [param0], %rd1;350; CHECK-NEXT: call.uni (retval0), test_s_i8f64p, (param0);351; CHECK-NEXT: ld.param.b64 %rd5, [retval0+24];352; CHECK-NEXT: ld.param.b8 %rs1, [retval0+8];353; CHECK-NEXT: ld.param.b64 %rd6, [retval0];354; CHECK-NEXT: ld.param.b8 %rd7, [retval0+16];355; CHECK-NEXT: ld.param.b8 %rd8, [retval0+15];356; CHECK-NEXT: ld.param.b8 %rd9, [retval0+14];357; CHECK-NEXT: ld.param.b8 %rd10, [retval0+13];358; CHECK-NEXT: ld.param.b8 %rd11, [retval0+12];359; CHECK-NEXT: ld.param.b8 %rd12, [retval0+11];360; CHECK-NEXT: ld.param.b8 %rd13, [retval0+10];361; CHECK-NEXT: ld.param.b8 %rd14, [retval0+9];362; CHECK-NEXT: } // callseq 6363; CHECK-NEXT: shl.b64 %rd15, %rd13, 8;364; CHECK-NEXT: or.b64 %rd16, %rd15, %rd14;365; CHECK-NEXT: shl.b64 %rd17, %rd12, 16;366; CHECK-NEXT: shl.b64 %rd18, %rd11, 24;367; CHECK-NEXT: or.b64 %rd19, %rd18, %rd17;368; CHECK-NEXT: or.b64 %rd20, %rd19, %rd16;369; CHECK-NEXT: shl.b64 %rd21, %rd9, 8;370; CHECK-NEXT: or.b64 %rd22, %rd21, %rd10;371; CHECK-NEXT: shl.b64 %rd23, %rd8, 16;372; CHECK-NEXT: shl.b64 %rd24, %rd7, 24;373; CHECK-NEXT: or.b64 %rd25, %rd24, %rd23;374; CHECK-NEXT: or.b64 %rd26, %rd25, %rd22;375; CHECK-NEXT: shl.b64 %rd27, %rd26, 32;376; CHECK-NEXT: or.b64 %rd28, %rd27, %rd20;377; CHECK-NEXT: st.param.b8 [func_retval0+9], %rd14;378; CHECK-NEXT: st.param.b64 [func_retval0+24], %rd5;379; CHECK-NEXT: st.param.b8 [func_retval0+8], %rs1;380; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;381; CHECK-NEXT: shr.u64 %rd29, %rd28, 56;382; CHECK-NEXT: st.param.b8 [func_retval0+16], %rd29;383; CHECK-NEXT: shr.u64 %rd30, %rd28, 48;384; CHECK-NEXT: st.param.b8 [func_retval0+15], %rd30;385; CHECK-NEXT: shr.u64 %rd31, %rd28, 40;386; CHECK-NEXT: st.param.b8 [func_retval0+14], %rd31;387; CHECK-NEXT: shr.u64 %rd32, %rd28, 32;388; CHECK-NEXT: st.param.b8 [func_retval0+13], %rd32;389; CHECK-NEXT: shr.u64 %rd33, %rd28, 24;390; CHECK-NEXT: st.param.b8 [func_retval0+12], %rd33;391; CHECK-NEXT: shr.u64 %rd34, %rd28, 16;392; CHECK-NEXT: st.param.b8 [func_retval0+11], %rd34;393; CHECK-NEXT: shr.u64 %rd35, %rd28, 8;394; CHECK-NEXT: st.param.b8 [func_retval0+10], %rd35;395; CHECK-NEXT: ret;396 %r = tail call %s_i8f64p @test_s_i8f64p(%s_i8f64p %a)397 ret %s_i8f64p %r398}399