[llvm] [ExpandVariadics] Fixup handeling of mismatched return types (PR #198843)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Fri May 22 14:26:03 PDT 2026
https://github.com/AlexMaclean updated https://github.com/llvm/llvm-project/pull/198843
>From cba655e76f4c2b365ce64129212a57379b4d3253 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Wed, 20 May 2026 10:07:31 -0700
Subject: [PATCH] [ExpandVariadics] Fixup handeling of mismatched return types
---
llvm/lib/Transforms/IPO/ExpandVariadics.cpp | 25 ++--
.../CodeGen/NVPTX/convert-call-to-indirect.ll | 128 ++++++++++--------
2 files changed, 90 insertions(+), 63 deletions(-)
diff --git a/llvm/lib/Transforms/IPO/ExpandVariadics.cpp b/llvm/lib/Transforms/IPO/ExpandVariadics.cpp
index 25715d267b242..5a0cce11fdef6 100644
--- a/llvm/lib/Transforms/IPO/ExpandVariadics.cpp
+++ b/llvm/lib/Transforms/IPO/ExpandVariadics.cpp
@@ -52,6 +52,7 @@
//===----------------------------------------------------------------------===//
#include "llvm/Transforms/IPO/ExpandVariadics.h"
+#include "llvm/ADT/Sequence.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/Demangle/Demangle.h"
#include "llvm/IR/IRBuilder.h"
@@ -232,12 +233,16 @@ class ExpandVariadics : public ModulePass {
bool expandVAIntrinsicCall(IRBuilder<> &Builder, const DataLayout &DL,
VACopyInst *Inst);
- FunctionType *inlinableVariadicFunctionType(Module &M, FunctionType *FTy) {
+ FunctionType *inlinableVariadicFunctionType(Module &M, FunctionType *FTy,
+ Type *ReturnType) {
// The type of "FTy" with the ... removed and a va_list appended
SmallVector<Type *> ArgTypes(FTy->params());
ArgTypes.push_back(ABI->vaListParameterType(M));
- return FunctionType::get(FTy->getReturnType(), ArgTypes,
- /*IsVarArgs=*/false);
+ return FunctionType::get(ReturnType, ArgTypes, /*IsVarArgs=*/false);
+ }
+
+ FunctionType *inlinableVariadicFunctionType(Module &M, FunctionType *FTy) {
+ return inlinableVariadicFunctionType(M, FTy, FTy->getReturnType());
}
bool expansionApplicableToFunction(Module &M, Function *F) {
@@ -645,12 +650,9 @@ bool ExpandVariadics::expandCall(Module &M, IRBuilder<> &Builder, CallBase *CB,
// This is tricky. The call instruction's function type might not match
// the type of the caller. When optimising, can leave it unchanged.
// Webassembly detects that inconsistency and repairs it.
- FunctionType *FuncType = CB->getFunctionType();
- if (FuncType != VarargFunctionType) {
+ if (CB->getFunctionType() != VarargFunctionType)
if (!rewriteABI())
return Changed;
- FuncType = VarargFunctionType;
- }
auto &Ctx = CB->getContext();
@@ -668,7 +670,7 @@ bool ExpandVariadics::expandCall(Module &M, IRBuilder<> &Builder, CallBase *CB,
uint64_t CurrentOffset = 0;
- for (unsigned I = FuncType->getNumParams(), E = CB->arg_size(); I < E; ++I) {
+ for (unsigned I : seq(VarargFunctionType->getNumParams(), CB->arg_size())) {
Value *ArgVal = CB->getArgOperand(I);
const bool IsByVal = CB->paramHasAttr(I, Attribute::ByVal);
const bool IsByRef = CB->paramHasAttr(I, Attribute::ByRef);
@@ -773,7 +775,7 @@ bool ExpandVariadics::expandCall(Module &M, IRBuilder<> &Builder, CallBase *CB,
Builder.CreateLifetimeStart(Alloced);
Frame.initializeStructAlloca(DL, Builder, Alloced, VarargsTy);
- const unsigned NumArgs = FuncType->getNumParams();
+ const unsigned NumArgs = VarargFunctionType->getNumParams();
SmallVector<Value *> Args(CB->arg_begin(), CB->arg_begin() + NumArgs);
// Initialize a va_list pointing to that struct and pass it as the last
@@ -809,7 +811,10 @@ bool ExpandVariadics::expandCall(Module &M, IRBuilder<> &Builder, CallBase *CB,
if (CallInst *CI = dyn_cast<CallInst>(CB)) {
Value *Dst = NF ? NF : CI->getCalledOperand();
- FunctionType *NFTy = inlinableVariadicFunctionType(M, VarargFunctionType);
+ // Use the type of the call site rather than the function type to ensure
+ // RAUW succeeds in the case of a mismatching return type.
+ FunctionType *NFTy =
+ inlinableVariadicFunctionType(M, VarargFunctionType, CB->getType());
NewCB = CallInst::Create(NFTy, Dst, Args, OpBundles, "", CI->getIterator());
diff --git a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
index 3304f18473e7e..d82a42ae5d56e 100644
--- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
@@ -1,6 +1,6 @@
; 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-sm_90 %{ llc < %s -march=nvptx64 -mcpu=sm_90 | %ptxas-verify -arch=sm_90 %}
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 | FileCheck %s
+; RUN: %if ptxas-sm_90 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 | %ptxas-verify -arch=sm_90 %}
%struct.64 = type <{ i64 }>
declare i64 @callee(ptr %p);
@@ -111,55 +111,63 @@ define i64 @test_param_count_mismatch(ptr %p) {
define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-LABEL: test_return_type_mismatch_variadic(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<32>;
+; CHECK-NEXT: .local .align 1 .b8 __local_depot3[1];
+; CHECK-NEXT: .reg .b64 %SP;
+; CHECK-NEXT: .reg .b64 %SPL;
+; CHECK-NEXT: .reg .b64 %rd<33>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.b64 %SPL, __local_depot3;
+; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_variadic_param_0];
; CHECK-NEXT: { // callseq 3, 0
; CHECK-NEXT: .param .b64 param0;
+; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .align 1 .b8 retval0[8];
+; CHECK-NEXT: add.u64 %rd2, %SP, 0;
+; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; CHECK-NEXT: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
-; CHECK-NEXT: mov.b64 %rd2, callee_variadic;
-; CHECK-NEXT: call (retval0), %rd2, (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: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
+; CHECK-NEXT: mov.b64 %rd3, callee_variadic;
+; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_3;
+; CHECK-NEXT: ld.param.b8 %rd4, [retval0+7];
+; CHECK-NEXT: ld.param.b8 %rd5, [retval0+6];
+; CHECK-NEXT: ld.param.b8 %rd6, [retval0+5];
+; CHECK-NEXT: ld.param.b8 %rd7, [retval0+4];
+; CHECK-NEXT: ld.param.b8 %rd8, [retval0+3];
+; CHECK-NEXT: ld.param.b8 %rd9, [retval0+2];
+; CHECK-NEXT: ld.param.b8 %rd10, [retval0+1];
+; CHECK-NEXT: ld.param.b8 %rd11, [retval0];
; CHECK-NEXT: } // callseq 3
-; CHECK-NEXT: shl.b64 %rd11, %rd9, 8;
-; CHECK-NEXT: or.b64 %rd12, %rd11, %rd10;
-; CHECK-NEXT: shl.b64 %rd13, %rd8, 16;
-; CHECK-NEXT: shl.b64 %rd14, %rd7, 24;
-; CHECK-NEXT: or.b64 %rd15, %rd14, %rd13;
-; CHECK-NEXT: or.b64 %rd16, %rd15, %rd12;
-; CHECK-NEXT: shl.b64 %rd17, %rd5, 8;
-; CHECK-NEXT: or.b64 %rd18, %rd17, %rd6;
-; CHECK-NEXT: shl.b64 %rd19, %rd4, 16;
-; CHECK-NEXT: shl.b64 %rd20, %rd3, 24;
-; CHECK-NEXT: or.b64 %rd21, %rd20, %rd19;
-; CHECK-NEXT: or.b64 %rd22, %rd21, %rd18;
-; CHECK-NEXT: shl.b64 %rd23, %rd22, 32;
-; CHECK-NEXT: or.b64 %rd24, %rd23, %rd16;
-; CHECK-NEXT: st.param.b8 [func_retval0], %rd10;
-; CHECK-NEXT: shr.u64 %rd25, %rd24, 56;
-; CHECK-NEXT: st.param.b8 [func_retval0+7], %rd25;
-; CHECK-NEXT: shr.u64 %rd26, %rd24, 48;
-; CHECK-NEXT: st.param.b8 [func_retval0+6], %rd26;
-; CHECK-NEXT: shr.u64 %rd27, %rd24, 40;
-; CHECK-NEXT: st.param.b8 [func_retval0+5], %rd27;
-; CHECK-NEXT: shr.u64 %rd28, %rd24, 32;
-; CHECK-NEXT: st.param.b8 [func_retval0+4], %rd28;
-; CHECK-NEXT: shr.u64 %rd29, %rd24, 24;
-; CHECK-NEXT: st.param.b8 [func_retval0+3], %rd29;
-; CHECK-NEXT: shr.u64 %rd30, %rd24, 16;
-; CHECK-NEXT: st.param.b8 [func_retval0+2], %rd30;
-; CHECK-NEXT: shr.u64 %rd31, %rd24, 8;
-; CHECK-NEXT: st.param.b8 [func_retval0+1], %rd31;
+; CHECK-NEXT: shl.b64 %rd12, %rd10, 8;
+; CHECK-NEXT: or.b64 %rd13, %rd12, %rd11;
+; CHECK-NEXT: shl.b64 %rd14, %rd9, 16;
+; CHECK-NEXT: shl.b64 %rd15, %rd8, 24;
+; CHECK-NEXT: or.b64 %rd16, %rd15, %rd14;
+; CHECK-NEXT: or.b64 %rd17, %rd16, %rd13;
+; CHECK-NEXT: shl.b64 %rd18, %rd6, 8;
+; CHECK-NEXT: or.b64 %rd19, %rd18, %rd7;
+; CHECK-NEXT: shl.b64 %rd20, %rd5, 16;
+; CHECK-NEXT: shl.b64 %rd21, %rd4, 24;
+; CHECK-NEXT: or.b64 %rd22, %rd21, %rd20;
+; CHECK-NEXT: or.b64 %rd23, %rd22, %rd19;
+; CHECK-NEXT: shl.b64 %rd24, %rd23, 32;
+; CHECK-NEXT: or.b64 %rd25, %rd24, %rd17;
+; CHECK-NEXT: st.param.b8 [func_retval0], %rd11;
+; CHECK-NEXT: shr.u64 %rd26, %rd25, 56;
+; CHECK-NEXT: st.param.b8 [func_retval0+7], %rd26;
+; CHECK-NEXT: shr.u64 %rd27, %rd25, 48;
+; CHECK-NEXT: st.param.b8 [func_retval0+6], %rd27;
+; CHECK-NEXT: shr.u64 %rd28, %rd25, 40;
+; CHECK-NEXT: st.param.b8 [func_retval0+5], %rd28;
+; CHECK-NEXT: shr.u64 %rd29, %rd25, 32;
+; CHECK-NEXT: st.param.b8 [func_retval0+4], %rd29;
+; CHECK-NEXT: shr.u64 %rd30, %rd25, 24;
+; CHECK-NEXT: st.param.b8 [func_retval0+3], %rd30;
+; CHECK-NEXT: shr.u64 %rd31, %rd25, 16;
+; CHECK-NEXT: st.param.b8 [func_retval0+2], %rd31;
+; CHECK-NEXT: shr.u64 %rd32, %rd25, 8;
+; CHECK-NEXT: st.param.b8 [func_retval0+1], %rd32;
; CHECK-NEXT: ret;
%ret = call %struct.64 (ptr, ...) @callee_variadic(ptr %p)
ret %struct.64 %ret
@@ -168,20 +176,27 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
define i64 @test_param_type_mismatch_variadic(ptr %p) {
; CHECK-LABEL: test_param_type_mismatch_variadic(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .local .align 8 .b8 __local_depot4[8];
+; CHECK-NEXT: .reg .b64 %SP;
+; CHECK-NEXT: .reg .b64 %SPL;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.b64 %SPL, __local_depot4;
+; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_type_mismatch_variadic_param_0];
+; CHECK-NEXT: st.b64 [%SP], 7;
; CHECK-NEXT: { // callseq 4, 0
-; CHECK-NEXT: .param .align 8 .b8 param1[8];
; CHECK-NEXT: .param .b64 param0;
+; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b64 retval0;
+; CHECK-NEXT: add.u64 %rd2, %SP, 0;
+; CHECK-NEXT: st.param.b64 [param1], %rd2;
; 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: ld.param.b64 %rd3, [retval0];
; CHECK-NEXT: } // callseq 4
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%ret = call i64 (ptr, ...) @callee_variadic(ptr %p, i64 7)
ret i64 %ret
@@ -190,20 +205,27 @@ define i64 @test_param_type_mismatch_variadic(ptr %p) {
define i64 @test_param_count_mismatch_variadic(ptr %p) {
; CHECK-LABEL: test_param_count_mismatch_variadic(
; CHECK: {
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .local .align 8 .b8 __local_depot5[8];
+; CHECK-NEXT: .reg .b64 %SP;
+; CHECK-NEXT: .reg .b64 %SPL;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.b64 %SPL, __local_depot5;
+; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_variadic_param_0];
+; CHECK-NEXT: st.b64 [%SP], 7;
; CHECK-NEXT: { // callseq 5, 0
-; CHECK-NEXT: .param .align 8 .b8 param1[8];
; CHECK-NEXT: .param .b64 param0;
+; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b64 retval0;
+; CHECK-NEXT: add.u64 %rd2, %SP, 0;
+; CHECK-NEXT: st.param.b64 [param1], %rd2;
; 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: ld.param.b64 %rd3, [retval0];
; CHECK-NEXT: } // callseq 5
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%ret = call i64 (ptr, ...) @callee_variadic(ptr %p, i64 7)
ret i64 %ret
More information about the llvm-commits
mailing list