[llvm] [NVPTX] Add support for empty type params and returns (PR #207057)

[Alex MacLean via llvm-commits]

================
@@ -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;
----------------
AlexMaclean wrote:

Fixed!

https://github.com/llvm/llvm-project/pull/207057