diff options
Diffstat (limited to 'llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll')
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll | 172 |
1 files changed, 162 insertions, 10 deletions
diff --git a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll index d1b478d..48209a8 100644 --- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll +++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll @@ -1,3 +1,4 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 | FileCheck %s ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_90 | %ptxas-verify -arch=sm_90 %} @@ -7,52 +8,203 @@ declare i64 @callee_variadic(ptr %p, ...); define %struct.64 @test_return_type_mismatch(ptr %p) { ; CHECK-LABEL: test_return_type_mismatch( -; CHECK: .param .align 1 .b8 retval0[8]; +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<40>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd2, [test_return_type_mismatch_param_0]; +; CHECK-NEXT: { // callseq 0, 0 +; CHECK-NEXT: .param .b64 param0; +; CHECK-NEXT: .param .align 1 .b8 retval0[8]; +; CHECK-NEXT: st.param.b64 [param0], %rd2; ; CHECK-NEXT: prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _); -; CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0), prototype_0; +; CHECK-NEXT: mov.b64 %rd1, callee; +; CHECK-NEXT: call (retval0), %rd1, (param0), prototype_0; +; CHECK-NEXT: ld.param.b8 %rd3, [retval0+7]; +; CHECK-NEXT: ld.param.b8 %rd4, [retval0+6]; +; CHECK-NEXT: ld.param.b8 %rd5, [retval0+5]; +; CHECK-NEXT: ld.param.b8 %rd6, [retval0+4]; +; CHECK-NEXT: ld.param.b8 %rd7, [retval0+3]; +; CHECK-NEXT: ld.param.b8 %rd8, [retval0+2]; +; CHECK-NEXT: ld.param.b8 %rd9, [retval0+1]; +; CHECK-NEXT: ld.param.b8 %rd10, [retval0]; +; CHECK-NEXT: } // callseq 0 +; CHECK-NEXT: shl.b64 %rd13, %rd9, 8; +; CHECK-NEXT: or.b64 %rd14, %rd13, %rd10; +; CHECK-NEXT: shl.b64 %rd16, %rd8, 16; +; CHECK-NEXT: shl.b64 %rd18, %rd7, 24; +; CHECK-NEXT: or.b64 %rd19, %rd18, %rd16; +; CHECK-NEXT: or.b64 %rd20, %rd19, %rd14; +; CHECK-NEXT: shl.b64 %rd23, %rd5, 8; +; CHECK-NEXT: or.b64 %rd24, %rd23, %rd6; +; CHECK-NEXT: shl.b64 %rd26, %rd4, 16; +; CHECK-NEXT: shl.b64 %rd28, %rd3, 24; +; CHECK-NEXT: or.b64 %rd29, %rd28, %rd26; +; CHECK-NEXT: or.b64 %rd30, %rd29, %rd24; +; CHECK-NEXT: shl.b64 %rd31, %rd30, 32; +; CHECK-NEXT: or.b64 %rd32, %rd31, %rd20; +; CHECK-NEXT: st.param.b8 [func_retval0], %rd10; +; CHECK-NEXT: shr.u64 %rd33, %rd32, 56; +; CHECK-NEXT: st.param.b8 [func_retval0+7], %rd33; +; CHECK-NEXT: shr.u64 %rd34, %rd32, 48; +; CHECK-NEXT: st.param.b8 [func_retval0+6], %rd34; +; CHECK-NEXT: shr.u64 %rd35, %rd32, 40; +; CHECK-NEXT: st.param.b8 [func_retval0+5], %rd35; +; CHECK-NEXT: shr.u64 %rd36, %rd32, 32; +; CHECK-NEXT: st.param.b8 [func_retval0+4], %rd36; +; CHECK-NEXT: shr.u64 %rd37, %rd32, 24; +; CHECK-NEXT: st.param.b8 [func_retval0+3], %rd37; +; CHECK-NEXT: shr.u64 %rd38, %rd32, 16; +; CHECK-NEXT: st.param.b8 [func_retval0+2], %rd38; +; CHECK-NEXT: shr.u64 %rd39, %rd32, 8; +; CHECK-NEXT: st.param.b8 [func_retval0+1], %rd39; +; CHECK-NEXT: ret; %ret = call %struct.64 @callee(ptr %p) ret %struct.64 %ret } define i64 @test_param_type_mismatch(ptr %p) { ; CHECK-LABEL: test_param_type_mismatch( -; CHECK: .param .b64 retval0; +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: { // callseq 1, 0 +; CHECK-NEXT: .param .b64 param0; +; CHECK-NEXT: .param .b64 retval0; ; CHECK-NEXT: prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _); -; CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0), prototype_1; +; CHECK-NEXT: st.param.b64 [param0], 7; +; CHECK-NEXT: mov.b64 %rd1, callee; +; CHECK-NEXT: call (retval0), %rd1, (param0), prototype_1; +; CHECK-NEXT: ld.param.b64 %rd2, [retval0]; +; CHECK-NEXT: } // callseq 1 +; CHECK-NEXT: st.param.b64 [func_retval0], %rd2; +; CHECK-NEXT: ret; %ret = call i64 @callee(i64 7) ret i64 %ret } define i64 @test_param_count_mismatch(ptr %p) { ; CHECK-LABEL: test_param_count_mismatch( -; CHECK: .param .b64 retval0; +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd2, [test_param_count_mismatch_param_0]; +; CHECK-NEXT: { // callseq 2, 0 +; CHECK-NEXT: .param .b64 param0; +; CHECK-NEXT: .param .b64 param1; +; CHECK-NEXT: .param .b64 retval0; +; CHECK-NEXT: st.param.b64 [param0], %rd2; ; CHECK-NEXT: prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _); -; CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0, param1), prototype_2; +; CHECK-NEXT: st.param.b64 [param1], 7; +; CHECK-NEXT: mov.b64 %rd1, callee; +; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_2; +; CHECK-NEXT: ld.param.b64 %rd3, [retval0]; +; CHECK-NEXT: } // callseq 2 +; CHECK-NEXT: st.param.b64 [func_retval0], %rd3; +; CHECK-NEXT: ret; %ret = call i64 @callee(ptr %p, i64 7) ret i64 %ret } define %struct.64 @test_return_type_mismatch_variadic(ptr %p) { ; CHECK-LABEL: test_return_type_mismatch_variadic( -; CHECK: .param .align 1 .b8 retval0[8]; +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<40>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd2, [test_return_type_mismatch_variadic_param_0]; +; CHECK-NEXT: { // callseq 3, 0 +; CHECK-NEXT: .param .b64 param0; +; CHECK-NEXT: .param .align 1 .b8 retval0[8]; +; CHECK-NEXT: st.param.b64 [param0], %rd2; ; CHECK-NEXT: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _); -; CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0), prototype_3; +; CHECK-NEXT: mov.b64 %rd1, callee_variadic; +; CHECK-NEXT: call (retval0), %rd1, (param0), prototype_3; +; CHECK-NEXT: ld.param.b8 %rd3, [retval0+7]; +; CHECK-NEXT: ld.param.b8 %rd4, [retval0+6]; +; CHECK-NEXT: ld.param.b8 %rd5, [retval0+5]; +; CHECK-NEXT: ld.param.b8 %rd6, [retval0+4]; +; CHECK-NEXT: ld.param.b8 %rd7, [retval0+3]; +; CHECK-NEXT: ld.param.b8 %rd8, [retval0+2]; +; CHECK-NEXT: ld.param.b8 %rd9, [retval0+1]; +; CHECK-NEXT: ld.param.b8 %rd10, [retval0]; +; CHECK-NEXT: } // callseq 3 +; CHECK-NEXT: shl.b64 %rd13, %rd9, 8; +; CHECK-NEXT: or.b64 %rd14, %rd13, %rd10; +; CHECK-NEXT: shl.b64 %rd16, %rd8, 16; +; CHECK-NEXT: shl.b64 %rd18, %rd7, 24; +; CHECK-NEXT: or.b64 %rd19, %rd18, %rd16; +; CHECK-NEXT: or.b64 %rd20, %rd19, %rd14; +; CHECK-NEXT: shl.b64 %rd23, %rd5, 8; +; CHECK-NEXT: or.b64 %rd24, %rd23, %rd6; +; CHECK-NEXT: shl.b64 %rd26, %rd4, 16; +; CHECK-NEXT: shl.b64 %rd28, %rd3, 24; +; CHECK-NEXT: or.b64 %rd29, %rd28, %rd26; +; CHECK-NEXT: or.b64 %rd30, %rd29, %rd24; +; CHECK-NEXT: shl.b64 %rd31, %rd30, 32; +; CHECK-NEXT: or.b64 %rd32, %rd31, %rd20; +; CHECK-NEXT: st.param.b8 [func_retval0], %rd10; +; CHECK-NEXT: shr.u64 %rd33, %rd32, 56; +; CHECK-NEXT: st.param.b8 [func_retval0+7], %rd33; +; CHECK-NEXT: shr.u64 %rd34, %rd32, 48; +; CHECK-NEXT: st.param.b8 [func_retval0+6], %rd34; +; CHECK-NEXT: shr.u64 %rd35, %rd32, 40; +; CHECK-NEXT: st.param.b8 [func_retval0+5], %rd35; +; CHECK-NEXT: shr.u64 %rd36, %rd32, 32; +; CHECK-NEXT: st.param.b8 [func_retval0+4], %rd36; +; CHECK-NEXT: shr.u64 %rd37, %rd32, 24; +; CHECK-NEXT: st.param.b8 [func_retval0+3], %rd37; +; CHECK-NEXT: shr.u64 %rd38, %rd32, 16; +; CHECK-NEXT: st.param.b8 [func_retval0+2], %rd38; +; CHECK-NEXT: shr.u64 %rd39, %rd32, 8; +; CHECK-NEXT: st.param.b8 [func_retval0+1], %rd39; +; CHECK-NEXT: ret; %ret = call %struct.64 (ptr, ...) @callee_variadic(ptr %p) ret %struct.64 %ret } define i64 @test_param_type_mismatch_variadic(ptr %p) { ; CHECK-LABEL: test_param_type_mismatch_variadic( -; CHECK: .param .b64 retval0; +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [test_param_type_mismatch_variadic_param_0]; +; CHECK-NEXT: { // callseq 4, 0 +; CHECK-NEXT: .param .align 8 .b8 param1[8]; +; CHECK-NEXT: .param .b64 param0; +; CHECK-NEXT: .param .b64 retval0; +; CHECK-NEXT: st.param.b64 [param0], %rd1; +; CHECK-NEXT: st.param.b64 [param1], 7; ; CHECK-NEXT: call.uni (retval0), callee_variadic, (param0, param1); +; CHECK-NEXT: ld.param.b64 %rd2, [retval0]; +; CHECK-NEXT: } // callseq 4 +; CHECK-NEXT: st.param.b64 [func_retval0], %rd2; +; CHECK-NEXT: ret; %ret = call i64 (ptr, ...) @callee_variadic(ptr %p, i64 7) ret i64 %ret } define i64 @test_param_count_mismatch_variadic(ptr %p) { ; CHECK-LABEL: test_param_count_mismatch_variadic( -; CHECK: .param .b64 retval0; +; CHECK: { +; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_variadic_param_0]; +; CHECK-NEXT: { // callseq 5, 0 +; CHECK-NEXT: .param .align 8 .b8 param1[8]; +; CHECK-NEXT: .param .b64 param0; +; CHECK-NEXT: .param .b64 retval0; +; CHECK-NEXT: st.param.b64 [param0], %rd1; +; CHECK-NEXT: st.param.b64 [param1], 7; ; CHECK-NEXT: call.uni (retval0), callee_variadic, (param0, param1); +; CHECK-NEXT: ld.param.b64 %rd2, [retval0]; +; CHECK-NEXT: } // callseq 5 +; CHECK-NEXT: st.param.b64 [func_retval0], %rd2; +; CHECK-NEXT: ret; %ret = call i64 (ptr, ...) @callee_variadic(ptr %p, i64 7) ret i64 %ret } |