'TYPE_EMPTY_P' vs. code offloading [PR120308]

We've got 'gcc/stor-layout.cc:finalize_type_size':

    /* Handle empty records as per the x86-64 psABI.  */
    TYPE_EMPTY_P (type) = targetm.calls.empty_record_p (type);

(Indeed x86_64 is still the only target to define 'TARGET_EMPTY_RECORD_P',
calling 'gcc/tree.cc-default_is_empty_record'.)

And so it happens that for an empty struct used in code offloaded from x86_64
host (but not powerpc64le host, for example), we get to see 'TYPE_EMPTY_P' in
offloading compilation (where the offload targets (currently?) don't use it
themselves, and therefore aren't prepared to handle it).

For nvptx offloading compilation, this causes wrong code generation:
'ptxas [...] error : Call has wrong number of parameters', as nvptx code
generation for function definition doesn't pay attention to this flag (say, in
'gcc/config/nvptx/nvptx.cc:pass_in_memory', or whereever else would be
appropriate to handle that), but the generic code 'gcc/calls.cc:expand_call'
via 'gcc/function.cc:aggregate_value_p' does pay attention to it, and we thus
get mismatching function definition vs. function call.

This issue apparently isn't a problem for GCN offloading, but I don't know if
that's by design or by accident.

Richard Biener:
> It looks like TYPE_EMPTY_P is only used during RTL expansion for ABI
> purposes, so computing it during layout_type is premature as shown here.
>
> I would suggest to simply re-compute it at offload stream-in time.

(For avoidance of doubt, the additions to 'gcc.target/nvptx/abi-struct-arg.c',
'gcc.target/nvptx/abi-struct-ret.c' are not dependent on the offload streaming
code changes, but are just to mirror the changes to
'libgomp.oacc-c-c++-common/abi-struct-1.c'.)

	PR lto/120308
	gcc/
	* lto-streamer-out.cc (hash_tree): Don't handle 'TYPE_EMPTY_P' for
	'lto_stream_offload_p'.
	* tree-streamer-in.cc (unpack_ts_type_common_value_fields):
	Likewise.
	* tree-streamer-out.cc (pack_ts_type_common_value_fields):
	Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c: Add empty
	structure testing.
	gcc/testsuite/
	* gcc.target/nvptx/abi-struct-arg.c: Add empty structure testing.
	* gcc.target/nvptx/abi-struct-ret.c: Likewise.

(cherry picked from commit 9063810c86beee6274d745b91d8fb43a81c9683e)

6 files changed, 61 insertions, 3 deletions

diff --git a/gcc/lto-streamer-out.cc b/gcc/lto-streamer-out.cc
index a055d12..8efda29 100644
--- a/gcc/lto-streamer-out.cc
+++ b/gcc/lto-streamer-out.cc
@@ -1376,7 +1376,8 @@ hash_tree (struct streamer_tree_cache_d *cache, hash_map<tree, hashval_t> *map,
       hstate.commit_flag ();
       hstate.add_int (TYPE_PRECISION_RAW (t));
       hstate.add_int (TYPE_ALIGN (t));
-      hstate.add_int (TYPE_EMPTY_P (t));
+      if (!lto_stream_offload_p)
+	hstate.add_int (TYPE_EMPTY_P (t));
     }
 
   if (CODE_CONTAINS_STRUCT (code, TS_TRANSLATION_UNIT_DECL))
diff --git a/gcc/testsuite/gcc.target/nvptx/abi-struct-arg.c b/gcc/testsuite/gcc.target/nvptx/abi-struct-arg.c
index 54ae651..c2cc4de 100644
--- a/gcc/testsuite/gcc.target/nvptx/abi-struct-arg.c
+++ b/gcc/testsuite/gcc.target/nvptx/abi-struct-arg.c
@@ -3,12 +3,16 @@
 
 /* Struct arg.  Passed via pointer.  */
 
+typedef struct {} empty;  /* See 'gcc/doc/extend.texi', "Empty Structures".  */
 typedef struct {char a;} one;
 typedef struct {short a;} two;
 typedef struct {int a;} four;
 typedef struct {long long a;} eight;
 typedef struct {int a, b[12];} big;
 
+/* { dg-final { scan-assembler-times ".extern .func dcl_aempty \\(.param.u64 %\[_a-z0-9\]*\\);" 1 } } */
+void dcl_aempty (empty);
+
 /* { dg-final { scan-assembler-times ".extern .func dcl_aone \\(.param.u64 %\[_a-z0-9\]*\\);" 1 } } */
 void dcl_aone (one);
 
@@ -28,6 +32,7 @@ void dcl_abig (big);
 
 void test_1 (void)
 {
+  dcl_aempty (({empty t; t;}));
   dcl_aone (M (one, 1));
   dcl_atwo (M (two, 2));
   dcl_afour (M (four, 3));
@@ -35,6 +40,11 @@ void test_1 (void)
   dcl_abig (M (big, 5));
 }
 
+/* { dg-final { scan-assembler-times ".visible .func dfn_aempty \\(.param.u64 %\[_a-z0-9\]*\\)(?:;|\[\r\n\]+\{)" 2 } } */
+void dfn_aempty (empty empty)
+{
+}
+
 /* { dg-final { scan-assembler-times ".visible .func dfn_aone \\(.param.u64 %\[_a-z0-9\]*\\)(?:;|\[\r\n\]+\{)" 2 } } */
 void dfn_aone (one one)
 {
diff --git a/gcc/testsuite/gcc.target/nvptx/abi-struct-ret.c b/gcc/testsuite/gcc.target/nvptx/abi-struct-ret.c
index d48a82d..13e5021 100644
--- a/gcc/testsuite/gcc.target/nvptx/abi-struct-ret.c
+++ b/gcc/testsuite/gcc.target/nvptx/abi-struct-ret.c
@@ -3,12 +3,16 @@
 
 /* Struct return.  Returned via pointer.  */
 
+typedef struct {} empty;  /* See 'gcc/doc/extend.texi', "Empty Structures".  */
 typedef struct {char a;} one;
 typedef struct {short a;} two;
 typedef struct {int a;} four;
 typedef struct {long long a;} eight;
 typedef struct {int a, b[12];} big;
 
+/* { dg-final { scan-assembler-times ".extern .func dcl_rempty \\(.param.u64 %\[_a-z0-9\]*\\);" 1 } } */
+empty dcl_rempty (void);
+
 /* { dg-final { scan-assembler-times ".extern .func dcl_rone \\(.param.u64 %\[_a-z0-9\]*\\);" 1 } } */
 one dcl_rone (void);
 
@@ -26,6 +30,7 @@ big dcl_rbig (void);
 
 void test_1 (void)
 {
+  dcl_rempty ();
   dcl_rone ();
   dcl_rtwo ();
   dcl_rfour ();
@@ -35,6 +40,12 @@ void test_1 (void)
 
 #define M(T, v) ({T t; t.a = v; t;})
 
+/* { dg-final { scan-assembler-times ".visible .func dfn_rempty \\(.param.u64 %\[_a-z0-9\]*\\)(?:;|\[\r\n\]+\{)" 2 } } */
+empty dfn_rempty (void)
+{
+  return ({empty t; t;});
+}
+
 /* { dg-final { scan-assembler-times ".visible .func dfn_rone \\(.param.u64 %\[_a-z0-9\]*\\)(?:;|\[\r\n\]+\{)" 2 } } */
 one dfn_rone (void)
 {
diff --git a/gcc/tree-streamer-in.cc b/gcc/tree-streamer-in.cc
index 215350f..e57e6c65 100644
--- a/gcc/tree-streamer-in.cc
+++ b/gcc/tree-streamer-in.cc
@@ -36,6 +36,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "asan.h"
 #include "opts.h"
 #include "stor-layout.h"
+#include "hooks.h" /* For 'hook_bool_const_tree_false'.  */
 
 
 /* Read a STRING_CST from the string table in DATA_IN using input
@@ -386,7 +387,16 @@ unpack_ts_type_common_value_fields (struct bitpack_d *bp, tree expr)
     TYPE_STRING_FLAG (expr) = (unsigned) bp_unpack_value (bp, 1);
   if (AGGREGATE_TYPE_P (expr))
     TYPE_TYPELESS_STORAGE (expr) = (unsigned) bp_unpack_value (bp, 1);
-  TYPE_EMPTY_P (expr) = (unsigned) bp_unpack_value (bp, 1);
+  if (!lto_stream_offload_p)
+    TYPE_EMPTY_P (expr) = (unsigned) bp_unpack_value (bp, 1);
+  else
+    {
+      /* All offload targets use the default ('false') 'TARGET_EMPTY_RECORD_P'.
+	 If that ever changes, we'll have to properly initialize 'TYPE_EMPTY_P'
+	 here, see 'stor-layout.cc:finalize_type_size' and PR120308.  */
+      gcc_assert (targetm.calls.empty_record_p == hook_bool_const_tree_false);
+      TYPE_EMPTY_P (expr) = 0;
+    }
   if (FUNC_OR_METHOD_TYPE_P (expr))
     TYPE_NO_NAMED_ARGS_STDARG_P (expr) = (unsigned) bp_unpack_value (bp, 1);
   if (RECORD_OR_UNION_TYPE_P (expr))
diff --git a/gcc/tree-streamer-out.cc b/gcc/tree-streamer-out.cc
index 3422725..4d008f7 100644
--- a/gcc/tree-streamer-out.cc
+++ b/gcc/tree-streamer-out.cc
@@ -372,7 +372,8 @@ pack_ts_type_common_value_fields (struct bitpack_d *bp, tree expr)
     bp_pack_value (bp, TYPE_STRING_FLAG (expr), 1);
   if (AGGREGATE_TYPE_P (expr))
     bp_pack_value (bp, TYPE_TYPELESS_STORAGE (expr), 1);
-  bp_pack_value (bp, TYPE_EMPTY_P (expr), 1);
+  if (!lto_stream_offload_p)
+    bp_pack_value (bp, TYPE_EMPTY_P (expr), 1);
   if (FUNC_OR_METHOD_TYPE_P (expr))
     bp_pack_value (bp, TYPE_NO_NAMED_ARGS_STDARG_P (expr), 1);
   if (RECORD_OR_UNION_TYPE_P (expr))
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c
index 379e9fd..8078655 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c
@@ -2,6 +2,7 @@
 
 /* See also '../libgomp.c-c++-common/target-abi-struct-1-O0.c'.  */
 
+typedef struct {} empty;  /* See 'gcc/doc/extend.texi', "Empty Structures".  */
 typedef struct {char a;} schar;
 typedef struct {short a;} sshort;
 typedef struct {int a;} sint;
@@ -12,6 +13,14 @@ typedef struct {int a, b[12];} sint_13;
 
 #define M(T) ({T t; t.a = sizeof t; t;})
 
+static __SIZE_TYPE__ empty_a;
+#pragma acc declare create(empty_a)
+#pragma acc routine
+static empty rempty(void)
+{
+  return ({empty t; empty_a = sizeof t; t;});
+}
+
 #pragma acc routine
 static schar rschar(void)
 {
@@ -43,6 +52,21 @@ static sint_13 rsint_13(void)
 }
 
 #pragma acc routine
+static void aempty(empty empty)
+{
+  (void) empty;
+
+  __SIZE_TYPE__ empty_a_exp;
+#ifndef __cplusplus
+  empty_a_exp = 0;
+#else
+  empty_a_exp = sizeof (char);
+#endif
+  if (empty_a != empty_a_exp)
+    __builtin_abort();
+}
+
+#pragma acc routine
 static void aschar(schar schar)
 {
   if (schar.a != sizeof (char))
@@ -85,6 +109,7 @@ int main()
 #pragma acc serial
   /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */
   {
+    aempty(rempty());
     aschar(rschar());
     asshort(rsshort());
     asint(rsint());