211 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 | FileCheck %s3; RUN: %if ptxas-sm_90 %{ llc < %s -march=nvptx64 -mcpu=sm_90 | %ptxas-verify -arch=sm_90 %}4 5%struct.64 = type <{ i64 }>6declare i64 @callee(ptr %p);7declare i64 @callee_variadic(ptr %p, ...);8 9define %struct.64 @test_return_type_mismatch(ptr %p) {10; CHECK-LABEL: test_return_type_mismatch(11; CHECK: {12; CHECK-NEXT: .reg .b64 %rd<32>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_param_0];16; CHECK-NEXT: { // callseq 0, 017; CHECK-NEXT: .param .b64 param0;18; CHECK-NEXT: .param .align 1 .b8 retval0[8];19; CHECK-NEXT: st.param.b64 [param0], %rd1;20; CHECK-NEXT: prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);21; CHECK-NEXT: mov.b64 %rd2, callee;22; CHECK-NEXT: call (retval0), %rd2, (param0), prototype_0;23; CHECK-NEXT: ld.param.b8 %rd3, [retval0+7];24; CHECK-NEXT: ld.param.b8 %rd4, [retval0+6];25; CHECK-NEXT: ld.param.b8 %rd5, [retval0+5];26; CHECK-NEXT: ld.param.b8 %rd6, [retval0+4];27; CHECK-NEXT: ld.param.b8 %rd7, [retval0+3];28; CHECK-NEXT: ld.param.b8 %rd8, [retval0+2];29; CHECK-NEXT: ld.param.b8 %rd9, [retval0+1];30; CHECK-NEXT: ld.param.b8 %rd10, [retval0];31; CHECK-NEXT: } // callseq 032; CHECK-NEXT: shl.b64 %rd11, %rd9, 8;33; CHECK-NEXT: or.b64 %rd12, %rd11, %rd10;34; CHECK-NEXT: shl.b64 %rd13, %rd8, 16;35; CHECK-NEXT: shl.b64 %rd14, %rd7, 24;36; CHECK-NEXT: or.b64 %rd15, %rd14, %rd13;37; CHECK-NEXT: or.b64 %rd16, %rd15, %rd12;38; CHECK-NEXT: shl.b64 %rd17, %rd5, 8;39; CHECK-NEXT: or.b64 %rd18, %rd17, %rd6;40; CHECK-NEXT: shl.b64 %rd19, %rd4, 16;41; CHECK-NEXT: shl.b64 %rd20, %rd3, 24;42; CHECK-NEXT: or.b64 %rd21, %rd20, %rd19;43; CHECK-NEXT: or.b64 %rd22, %rd21, %rd18;44; CHECK-NEXT: shl.b64 %rd23, %rd22, 32;45; CHECK-NEXT: or.b64 %rd24, %rd23, %rd16;46; CHECK-NEXT: st.param.b8 [func_retval0], %rd10;47; CHECK-NEXT: shr.u64 %rd25, %rd24, 56;48; CHECK-NEXT: st.param.b8 [func_retval0+7], %rd25;49; CHECK-NEXT: shr.u64 %rd26, %rd24, 48;50; CHECK-NEXT: st.param.b8 [func_retval0+6], %rd26;51; CHECK-NEXT: shr.u64 %rd27, %rd24, 40;52; CHECK-NEXT: st.param.b8 [func_retval0+5], %rd27;53; CHECK-NEXT: shr.u64 %rd28, %rd24, 32;54; CHECK-NEXT: st.param.b8 [func_retval0+4], %rd28;55; CHECK-NEXT: shr.u64 %rd29, %rd24, 24;56; CHECK-NEXT: st.param.b8 [func_retval0+3], %rd29;57; CHECK-NEXT: shr.u64 %rd30, %rd24, 16;58; CHECK-NEXT: st.param.b8 [func_retval0+2], %rd30;59; CHECK-NEXT: shr.u64 %rd31, %rd24, 8;60; CHECK-NEXT: st.param.b8 [func_retval0+1], %rd31;61; CHECK-NEXT: ret;62 %ret = call %struct.64 @callee(ptr %p)63 ret %struct.64 %ret64}65 66define i64 @test_param_type_mismatch(ptr %p) {67; CHECK-LABEL: test_param_type_mismatch(68; CHECK: {69; CHECK-NEXT: .reg .b64 %rd<3>;70; CHECK-EMPTY:71; CHECK-NEXT: // %bb.0:72; CHECK-NEXT: { // callseq 1, 073; CHECK-NEXT: .param .b64 param0;74; CHECK-NEXT: .param .b64 retval0;75; CHECK-NEXT: prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);76; CHECK-NEXT: st.param.b64 [param0], 7;77; CHECK-NEXT: mov.b64 %rd1, callee;78; CHECK-NEXT: call (retval0), %rd1, (param0), prototype_1;79; CHECK-NEXT: ld.param.b64 %rd2, [retval0];80; CHECK-NEXT: } // callseq 181; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;82; CHECK-NEXT: ret;83 %ret = call i64 @callee(i64 7)84 ret i64 %ret85}86 87define i64 @test_param_count_mismatch(ptr %p) {88; CHECK-LABEL: test_param_count_mismatch(89; CHECK: {90; CHECK-NEXT: .reg .b64 %rd<4>;91; CHECK-EMPTY:92; CHECK-NEXT: // %bb.0:93; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_param_0];94; CHECK-NEXT: { // callseq 2, 095; CHECK-NEXT: .param .b64 param0;96; CHECK-NEXT: .param .b64 param1;97; CHECK-NEXT: .param .b64 retval0;98; CHECK-NEXT: st.param.b64 [param0], %rd1;99; CHECK-NEXT: prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);100; CHECK-NEXT: st.param.b64 [param1], 7;101; CHECK-NEXT: mov.b64 %rd2, callee;102; CHECK-NEXT: call (retval0), %rd2, (param0, param1), prototype_2;103; CHECK-NEXT: ld.param.b64 %rd3, [retval0];104; CHECK-NEXT: } // callseq 2105; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;106; CHECK-NEXT: ret;107 %ret = call i64 @callee(ptr %p, i64 7)108 ret i64 %ret109}110 111define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {112; CHECK-LABEL: test_return_type_mismatch_variadic(113; CHECK: {114; CHECK-NEXT: .reg .b64 %rd<32>;115; CHECK-EMPTY:116; CHECK-NEXT: // %bb.0:117; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_variadic_param_0];118; CHECK-NEXT: { // callseq 3, 0119; CHECK-NEXT: .param .b64 param0;120; CHECK-NEXT: .param .align 1 .b8 retval0[8];121; CHECK-NEXT: st.param.b64 [param0], %rd1;122; CHECK-NEXT: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);123; CHECK-NEXT: mov.b64 %rd2, callee_variadic;124; CHECK-NEXT: call (retval0), %rd2, (param0), prototype_3;125; CHECK-NEXT: ld.param.b8 %rd3, [retval0+7];126; CHECK-NEXT: ld.param.b8 %rd4, [retval0+6];127; CHECK-NEXT: ld.param.b8 %rd5, [retval0+5];128; CHECK-NEXT: ld.param.b8 %rd6, [retval0+4];129; CHECK-NEXT: ld.param.b8 %rd7, [retval0+3];130; CHECK-NEXT: ld.param.b8 %rd8, [retval0+2];131; CHECK-NEXT: ld.param.b8 %rd9, [retval0+1];132; CHECK-NEXT: ld.param.b8 %rd10, [retval0];133; CHECK-NEXT: } // callseq 3134; CHECK-NEXT: shl.b64 %rd11, %rd9, 8;135; CHECK-NEXT: or.b64 %rd12, %rd11, %rd10;136; CHECK-NEXT: shl.b64 %rd13, %rd8, 16;137; CHECK-NEXT: shl.b64 %rd14, %rd7, 24;138; CHECK-NEXT: or.b64 %rd15, %rd14, %rd13;139; CHECK-NEXT: or.b64 %rd16, %rd15, %rd12;140; CHECK-NEXT: shl.b64 %rd17, %rd5, 8;141; CHECK-NEXT: or.b64 %rd18, %rd17, %rd6;142; CHECK-NEXT: shl.b64 %rd19, %rd4, 16;143; CHECK-NEXT: shl.b64 %rd20, %rd3, 24;144; CHECK-NEXT: or.b64 %rd21, %rd20, %rd19;145; CHECK-NEXT: or.b64 %rd22, %rd21, %rd18;146; CHECK-NEXT: shl.b64 %rd23, %rd22, 32;147; CHECK-NEXT: or.b64 %rd24, %rd23, %rd16;148; CHECK-NEXT: st.param.b8 [func_retval0], %rd10;149; CHECK-NEXT: shr.u64 %rd25, %rd24, 56;150; CHECK-NEXT: st.param.b8 [func_retval0+7], %rd25;151; CHECK-NEXT: shr.u64 %rd26, %rd24, 48;152; CHECK-NEXT: st.param.b8 [func_retval0+6], %rd26;153; CHECK-NEXT: shr.u64 %rd27, %rd24, 40;154; CHECK-NEXT: st.param.b8 [func_retval0+5], %rd27;155; CHECK-NEXT: shr.u64 %rd28, %rd24, 32;156; CHECK-NEXT: st.param.b8 [func_retval0+4], %rd28;157; CHECK-NEXT: shr.u64 %rd29, %rd24, 24;158; CHECK-NEXT: st.param.b8 [func_retval0+3], %rd29;159; CHECK-NEXT: shr.u64 %rd30, %rd24, 16;160; CHECK-NEXT: st.param.b8 [func_retval0+2], %rd30;161; CHECK-NEXT: shr.u64 %rd31, %rd24, 8;162; CHECK-NEXT: st.param.b8 [func_retval0+1], %rd31;163; CHECK-NEXT: ret;164 %ret = call %struct.64 (ptr, ...) @callee_variadic(ptr %p)165 ret %struct.64 %ret166}167 168define i64 @test_param_type_mismatch_variadic(ptr %p) {169; CHECK-LABEL: test_param_type_mismatch_variadic(170; CHECK: {171; CHECK-NEXT: .reg .b64 %rd<3>;172; CHECK-EMPTY:173; CHECK-NEXT: // %bb.0:174; CHECK-NEXT: ld.param.b64 %rd1, [test_param_type_mismatch_variadic_param_0];175; CHECK-NEXT: { // callseq 4, 0176; CHECK-NEXT: .param .align 8 .b8 param1[8];177; CHECK-NEXT: .param .b64 param0;178; CHECK-NEXT: .param .b64 retval0;179; CHECK-NEXT: st.param.b64 [param0], %rd1;180; CHECK-NEXT: st.param.b64 [param1], 7;181; CHECK-NEXT: call.uni (retval0), callee_variadic, (param0, param1);182; CHECK-NEXT: ld.param.b64 %rd2, [retval0];183; CHECK-NEXT: } // callseq 4184; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;185; CHECK-NEXT: ret;186 %ret = call i64 (ptr, ...) @callee_variadic(ptr %p, i64 7)187 ret i64 %ret188}189 190define i64 @test_param_count_mismatch_variadic(ptr %p) {191; CHECK-LABEL: test_param_count_mismatch_variadic(192; CHECK: {193; CHECK-NEXT: .reg .b64 %rd<3>;194; CHECK-EMPTY:195; CHECK-NEXT: // %bb.0:196; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_variadic_param_0];197; CHECK-NEXT: { // callseq 5, 0198; CHECK-NEXT: .param .align 8 .b8 param1[8];199; CHECK-NEXT: .param .b64 param0;200; CHECK-NEXT: .param .b64 retval0;201; CHECK-NEXT: st.param.b64 [param0], %rd1;202; CHECK-NEXT: st.param.b64 [param1], 7;203; CHECK-NEXT: call.uni (retval0), callee_variadic, (param0, param1);204; CHECK-NEXT: ld.param.b64 %rd2, [retval0];205; CHECK-NEXT: } // callseq 5206; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;207; CHECK-NEXT: ret;208 %ret = call i64 (ptr, ...) @callee_variadic(ptr %p, i64 7)209 ret i64 %ret210}211