From a835f046cdf017b9e8ad5576df4f10daaf8420d0 Mon Sep 17 00:00:00 2001
From: Tobias Burnus <tobias@codesourcery.com>
Date: Fri, 12 May 2023 16:27:40 +0200
Subject: [PATCH] LTO: Fix writing of toplevel asm with offloading [PR109816]

When offloading was enabled, top-level 'asm' were added to the offloading
section, confusing assemblers which did not support the syntax. Additionally,
with offloading and -flto, the top-level assembler code did not end up
in the host files.

As r14-321-g9a41d2cdbcd added top-level 'asm' to one libstdc++ header file,
the issue became more apparent, causing fails with nvptx for some
C++ testcases.

	PR libstdc++/109816

gcc/ChangeLog:

	* lto-cgraph.cc (output_symtab): Guard lto_output_toplevel_asms by
	'!lto_stream_offload_p'.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/target-map-class-1.C: New test.
	* testsuite/libgomp.c++/target-map-class-2.C: New test.
---
 gcc/lto-cgraph.cc                             |  2 +-
 .../libgomp.c++/target-map-class-1.C          | 98 +++++++++++++++++++
 .../libgomp.c++/target-map-class-2.C          |  6 ++
 3 files changed, 105 insertions(+), 1 deletion(-)
 create mode 100644 libgomp/testsuite/libgomp.c++/target-map-class-1.C
 create mode 100644 libgomp/testsuite/libgomp.c++/target-map-class-2.C

diff --git a/gcc/lto-cgraph.cc b/gcc/lto-cgraph.cc
index 805c7855eb3..aed5e9ddb18 100644
--- a/gcc/lto-cgraph.cc
+++ b/gcc/lto-cgraph.cc
@@ -1020,7 +1020,7 @@ output_symtab (void)
      When doing WPA we must output every asm just once.  Since we do not partition asm
      nodes at all, output them to first output.  This is kind of hack, but should work
      well.  */
-  if (!asm_nodes_output)
+  if (!asm_nodes_output && !lto_stream_offload_p)
     {
       asm_nodes_output = true;
       lto_output_toplevel_asms ();
diff --git a/libgomp/testsuite/libgomp.c++/target-map-class-1.C b/libgomp/testsuite/libgomp.c++/target-map-class-1.C
new file mode 100644
index 00000000000..ad4802d2766
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-map-class-1.C
@@ -0,0 +1,98 @@
+/* PR middle-end/109816  */
+
+/* This variant: without -flto, see target-map-class-2.C for -flto. */
+
+/* iostream.h adds 'globl _ZSt21ios_base_library_initv' with _GLIBCXX_SYMVER_GNU,
+   but it shouldn't end up in the offload assembly but only in the host assembly. */
+
+/* Example based on sollve_vv's test_target_data_map_classes.cpp; however,
+   relevant is only the 'include' and not the actual executable code.  */
+
+#include <iostream>
+#include <omp.h>
+
+using namespace std;
+
+#define N 1000
+
+struct A
+{
+  int *h_array;
+  int size, sum;
+
+  A (int *array, const int s) : h_array(array), size(s), sum(0) { }
+  ~A() { h_array = NULL; }
+};
+
+void
+test_map_tofrom_class_heap ()
+{
+  int *array = new int[N];
+  A *obj = new A (array, N);
+
+  #pragma omp target map(from: array[:N]) map(tofrom: obj[:1])
+    {
+      int *tmp_h_array = obj->h_array;
+      obj->h_array = array;
+      int tmp = 0;
+      for (int i = 0; i < N; ++i)
+	{
+	  obj->h_array[i] = 4*i;
+	  tmp += 3;
+	}
+      obj->h_array = tmp_h_array;
+      obj->sum = tmp;
+    }
+
+  for (int i = 0; i < N; ++i)
+    if (obj->h_array[i] != 4*i)
+      __builtin_abort ();
+
+  if (3*N != obj->sum)
+    {
+      std::cout << "sum: " << obj->sum << std::endl;
+      __builtin_abort ();
+    }
+
+  delete obj;
+  delete[] array;
+}
+
+void
+test_map_tofrom_class_stack ()
+{
+  int array[N];
+  A obj(array, N);
+
+  #pragma omp target map(from: array[:N]) map(tofrom: obj)
+    {
+      int *tmp_h_array = obj.h_array;
+      obj.h_array = array;
+      int tmp = 0;
+      for (int i = 0; i < N; ++i)
+	{
+	  obj.h_array[i] = 7*i;
+	  tmp += 5;
+	}
+      obj.h_array = tmp_h_array;
+      obj.sum = tmp;
+    }
+
+  for (int i = 0; i < N; ++i)
+    if (obj.h_array[i] != 7*i)
+      __builtin_abort ();
+
+  if (5*N != obj.sum)
+    {
+      std::cout << "sum: " << obj.sum << std::endl;
+      __builtin_abort ();
+    }
+}
+
+int
+main()
+{
+  test_map_tofrom_class_heap();
+  test_map_tofrom_class_stack();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-map-class-2.C b/libgomp/testsuite/libgomp.c++/target-map-class-2.C
new file mode 100644
index 00000000000..1ef20f7ec22
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-map-class-2.C
@@ -0,0 +1,6 @@
+/* { dg-additional-options "-flto" }  */
+/* PR middle-end/109816  */
+
+/* This variant: with -flto, see target-map-class-1.C for without -flto. */
+
+#include "target-map-class-1.C"