[llvm] [NVPTX] Add support for empty type params and returns (PR #207057)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 1 11:57:13 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Alex MacLean (AlexMaclean)
<details>
<summary>Changes</summary>
---
Full diff: https://github.com/llvm/llvm-project/pull/207057.diff
3 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (+17-4)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+5-3)
- (modified) llvm/test/CodeGen/NVPTX/empty-type.ll (+144-5)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 94acde4b22b5e..cbef40e49d9bb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -284,7 +284,9 @@ void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
Type *Ty = F->getReturnType();
- if (Ty->getTypeID() == Type::VoidTyID)
+ // A void or zero-sized return type (e.g. an empty struct) produces no return
+ // parameter.
+ if (Ty->isVoidTy() || Ty->isEmptyTy())
return;
O << " (";
@@ -325,7 +327,7 @@ void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
O << "prototype_" << UniqueCallSite << " : .callprototype ";
- if (RetTy->isVoidTy()) {
+ if (RetTy->isVoidTy() || RetTy->isEmptyTy()) {
O << "()";
} else {
O << "(";
@@ -403,10 +405,16 @@ void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
const FunctionType *FTy = CB.getFunctionType();
const unsigned NumArgs = FTy->getNumParams();
- interleave(seq(NumArgs), O, MakeArg, ", ");
+ // Zero-sized arguments (e.g. empty structs) are not passed and so do not
+ // appear in the prototype.
+ auto LiveArgs = make_filter_range(seq(NumArgs), [&](unsigned I) {
+ return !CB.getArgOperand(I)->getType()->isEmptyTy();
+ });
+
+ interleave(LiveArgs, O, MakeArg, ", ");
if (FTy->isVarArg() && CB.arg_size() > NumArgs)
- O << (NumArgs ? "," : "") << " .param .align "
+ O << (LiveArgs.empty() ? "" : ",") << " .param .align "
<< STI.getMaxRequiredAlignment() << " .b8 _[]";
O << ")";
@@ -1472,6 +1480,11 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
for (const Argument &Arg : F->args()) {
Type *Ty = Arg.getType();
+
+ // Skip zero-sized arguments; they are not passed as PTX parameters.
+ if (Ty->isEmptyTy())
+ continue;
+
const std::string ParamSym = TLI->getParamName(F, Arg.getArgNo());
if (!IsFirst)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 4047738a2d447..92ff303d9c1a2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1320,7 +1320,7 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
//
// After all vararg is processed, 'VAOffset' holds the size of the
// vararg byte array.
- assert((CLI.IsVarArg || CLI.Args.size() == CLI.NumFixedArgs) &&
+ assert((CLI.IsVarArg || CLI.Args.size() <= CLI.NumFixedArgs) &&
"Non-VarArg function with extra arguments");
const unsigned FirstVAArg = CLI.NumFixedArgs; // position of first variadic
@@ -3972,8 +3972,10 @@ SDValue NVPTXTargetLowering::LowerFormalArguments(
Type *Ty = Arg.getType();
- if (ArgIns.empty())
- report_fatal_error("Empty parameter types are not supported");
+ if (ArgIns.empty()) {
+ assert(Ty->isEmptyTy() && "Only empty types produce no parameter values");
+ continue;
+ }
if (Arg.use_empty()) {
// argument is dead
diff --git a/llvm/test/CodeGen/NVPTX/empty-type.ll b/llvm/test/CodeGen/NVPTX/empty-type.ll
index 7e84f109cf8b1..b384eb5d0ae7d 100644
--- a/llvm/test/CodeGen/NVPTX/empty-type.ll
+++ b/llvm/test/CodeGen/NVPTX/empty-type.ll
@@ -1,10 +1,149 @@
-; RUN: not --crash llc < %s -mtriple=nvptx -mcpu=sm_20 2>&1 | FileCheck %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 | FileCheck %s
+; RUN: %if ptxas-sm_90 %{ llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 | %ptxas-verify -arch=sm_90 %}
-%struct.A = type { [0 x float] }
-%struct.B = type { i32, i32 }
+; Zero-sized types (empty structs, zero-length arrays and aggregates built only
+; from them) do not decompose into any PTX values. No parameter or return value
+; is emitted for them, and the parameters that remain are numbered contiguously.
-; CHECK: ERROR: Empty parameter types are not supported
-define void @kernel(%struct.A %a, %struct.B %b) {
+%empty = type {}
+%zero_array = type { [0 x float] }
+%nested = type { {}, [0 x float] }
+
+; An empty struct return type is lowered like a void return.
+define %empty @ret_empty(i32 %n) {
+; CHECK-LABEL: ret_empty(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ret;
+ ret %empty zeroinitializer
+}
+
+; An empty struct as the only argument produces an empty parameter list.
+define void @only_empty(%empty %z) {
+; CHECK-LABEL: only_empty(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ret;
+ ret void
+}
+
+; A zero-sized argument between two real ones is dropped; the parameters that
+; remain keep their original argument-numbered names in the callee.
+define i32 @middle(i32 %a, %empty %z, i32 %b) {
+; CHECK-LABEL: middle(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [middle_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [middle_param_2];
+; CHECK-NEXT: add.s32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %s = add i32 %a, %b
+ ret i32 %s
+}
+
+; A zero-length array argument is dropped just like an empty struct.
+define i32 @zero_array(%zero_array %z, i32 %n) {
+; CHECK-LABEL: zero_array(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [zero_array_param_1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ ret i32 %n
+}
+
+; A nested aggregate that only contains zero-sized fields is also dropped.
+define i32 @nested(%nested %z, i32 %n) {
+; CHECK-LABEL: nested(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [nested_param_1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ ret i32 %n
+}
+
+; A kernel entry point with a zero-sized parameter.
+define ptx_kernel void @kernel(%empty %z, ptr %p, i32 %n) {
+; CHECK-LABEL: kernel(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [kernel_param_1];
+; CHECK-NEXT: cvta.to.global.u64 %rd2, %rd1;
+; CHECK-NEXT: ld.param.b32 %r1, [kernel_param_2];
+; CHECK-NEXT: st.global.b32 [%rd2], %r1;
+; CHECK-NEXT: ret;
+ store i32 %n, ptr %p
+ ret void
+}
+
+; A caller that passes zero-sized values and captures a zero-sized return value,
+; then forwards it on. No parameter or return slot is declared for them and the
+; passed parameters are numbered contiguously from param0.
+define void @caller() {
+; CHECK-LABEL: caller(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0: // %entry
+; CHECK-NEXT: { // callseq 0, 0
+; CHECK-NEXT: .param .b32 param0;
+; CHECK-NEXT: st.param.b32 [param0], 3;
+; CHECK-NEXT: call.uni ret_empty, (param0);
+; CHECK-NEXT: } // callseq 0
+; CHECK-NEXT: { // callseq 1, 0
+; CHECK-NEXT: call.uni only_empty, ();
+; CHECK-NEXT: } // callseq 1
+; CHECK-NEXT: { // callseq 2, 0
+; CHECK-NEXT: .param .b32 param0;
+; CHECK-NEXT: .param .b32 param1;
+; CHECK-NEXT: .param .b32 retval0;
+; CHECK-NEXT: st.param.b32 [param1], 2;
+; CHECK-NEXT: st.param.b32 [param0], 1;
+; CHECK-NEXT: call.uni (retval0), middle, (param0, param1);
+; CHECK-NEXT: } // callseq 2
+; CHECK-NEXT: { // callseq 3, 0
+; CHECK-NEXT: call.uni only_empty, ();
+; CHECK-NEXT: } // callseq 3
+; CHECK-NEXT: ret;
entry:
+ %r = call %empty @ret_empty(i32 3)
+ call void @only_empty(%empty zeroinitializer)
+ %m = call i32 @middle(i32 1, %empty zeroinitializer, i32 2)
+ call void @only_empty(%empty %r)
+ ret void
+}
+
+; An indirect call emits a matching prototype that also drops the zero-sized
+; argument and the empty return type.
+define void @indirect_caller(ptr %fp) {
+; CHECK-LABEL: indirect_caller(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-NEXT: prototype_4 : .callprototype ()_ (.param .b32 _);
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [indirect_caller_param_0];
+; CHECK-NEXT: { // callseq 4, 0
+; CHECK-NEXT: .param .b32 param0;
+; CHECK-NEXT: st.param.b32 [param0], 42;
+; CHECK-NEXT: call %rd1, (param0), prototype_4;
+; CHECK-NEXT: } // callseq 4
+; CHECK-NEXT: ret;
+ call void %fp(%empty zeroinitializer, i32 42)
ret void
}
``````````
</details>
https://github.com/llvm/llvm-project/pull/207057
More information about the llvm-commits
mailing list