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.
---
gcc/lto-streamer-out.cc | 3 ++-
.../gcc.target/nvptx/abi-struct-arg.c | 10 ++++++++
.../gcc.target/nvptx/abi-struct-ret.c | 11 ++++++++
gcc/tree-streamer-in.cc | 12 ++++++++-
gcc/tree-streamer-out.cc | 3 ++-
.../libgomp.oacc-c-c++-common/abi-struct-1.c | 25 +++++++++++++++++++
6 files changed, 61 insertions(+), 3 deletions(-)
diff --git a/gcc/lto-streamer-out.cc b/gcc/lto-streamer-out.cc
index 86d338461c0..308ab3416b3 100644
--- a/gcc/lto-streamer-out.cc
+++ b/gcc/lto-streamer-out.cc
@@ -1398,7 +1398,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 54ae651dcca..c2cc4de115e 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 d48a82d26ce..13e50212dc3 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 215350f6513..e57e6c6514a 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 34227259b8a..4d008f735ff 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 379e9fd3a97..80786555fe2 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)
{
@@ -42,6 +51,21 @@ static sint_13 rsint_13(void)
return M(sint_13);
}
+#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)
{
@@ -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());
--
2.34.1
