OpenMP: Constructors and destructors for "declare target" static aggregates

This commit also compile-time expands (__builtin_)omp_is_initial_device for
both Fortran and C/C++ (unless, -fno-builtin-omp_is_initial_device is used).
But the main change is:

This commit adds support for running constructors and destructors for
static (file-scope) aggregates for C++ objects which are marked with
"declare target" directives on OpenMP offload targets.

Before this commit, space is allocated on the target for such aggregates,
but nothing ever constructs them properly, so they end up zero-initialised.

(See the new test static-aggr-constructor-destructor-3.C for a reason
why running constructors on the target is preferable to e.g. constructing
on the host and then copying the resulting object to the target.)

2024-08-07  Julian Brown  <julian@codesourcery.com>
	    Tobias Burnus  <tobias@baylibre.com>

gcc/ChangeLog:

	* builtins.def (DEF_GOMP_BUILTIN_COMPILER): Define
	DEF_GOMP_BUILTIN_COMPILER to handle the non-prefix version.
	* gimple-fold.cc (gimple_fold_builtin_omp_is_initial_device): New.
	(gimple_fold_builtin): Call it.
	* omp-builtins.def (BUILT_IN_OMP_IS_INITIAL_DEVICE): Define.
	* tree.cc (get_file_function_name): Support names for on-target
	constructor/destructor functions.

gcc/cp/
	* decl2.cc (tree-inline.h): Include.
	(static_init_fini_fns): Bump to four entries. Update comment.
	(start_objects, start_partial_init_fini_fn): Add 'omp_target'
	parameter. Support "declare target" decls. Update forward declaration.
	(emit_partial_init_fini_fn): Add 'host_fn' parameter. Return tree for
	the created function. Support "declare target".
	(OMP_SSDF_IDENTIFIER): New macro.
	(partition_vars_for_init_fini): Support partitioning "declare target"
	variables also.
	(generate_ctor_or_dtor_function): Add 'omp_target' parameter. Support
	"declare target" decls.
	(c_parse_final_cleanups): Support constructors/destructors on OpenMP
	offload targets.

gcc/fortran/ChangeLog:

	* gfortran.h (gfc_option_t): Add disable_omp_is_initial_device.
	* lang.opt (fbuiltin-): Add.
	* options.cc (gfc_handle_option): Handle
	-fno-builtin-omp_is_initial_device.
	* f95-lang.cc (gfc_init_builtin_functions): Handle
	DEF_GOMP_BUILTIN_COMPILER.
	* trans-decl.cc (gfc_get_extern_function_decl): Add code to use
	DEF_GOMP_BUILTIN_COMPILER for 'omp_is_initial_device'.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test.
	* testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test.
	* testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test.
	* testsuite/libgomp.c-c++-common/target-is-initial-host.c: New test.
	* testsuite/libgomp.c-c++-common/target-is-initial-host-2.c: New test.
	* testsuite/libgomp.fortran/target-is-initial-host.f: New test.
	* testsuite/libgomp.fortran/target-is-initial-host.f90: New test.
	* testsuite/libgomp.fortran/target-is-initial-host-2.f90: New test.

Co-authored-by: Tobias Burnus <tobias@baylibre.com>

8 files changed, 350 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
new file mode 100644
index 0000000..403a071
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C
@@ -0,0 +1,72 @@
+// { dg-do run }
+// { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+// { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+// { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "_GLOBAL__off_I_v1" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "__omp_target_static_init_and_destruction" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 2 "gimple" } }
+
+// { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+// { dg-final { scan-tree-dump-not "__omp_target_static_init_and_destruction" "optimized" } }
+// FIXME: should be '-not' not '-times' 1:
+// { dg-final { scan-tree-dump-times "void _GLOBAL__off_I_v1" 1 "optimized" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 1 "optimized" } }
+
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_nvptx } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_nvptx } } }
+
+#include <cassert>
+#include <omp.h>
+
+#pragma omp declare target
+
+struct str {
+  str(int x) : _x(x) { }
+  int add(str o) { return _x + o._x; }
+  int _x;
+} v1(5);
+
+#pragma omp end declare target
+
+void check_host()
+{
+  assert (v1._x == 5);
+}
+
+void check_devs()
+{
+  for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+    {
+      int res = 99, dev_num = 98;
+      #pragma omp target map(from: res, dev_num) device(dev)
+	{
+	  res = v1._x;
+	  dev_num = omp_get_device_num();
+	}
+      assert (res == 5);
+      if (dev == omp_initial_device)
+	assert (dev_num == omp_get_num_devices());
+      else
+	assert (dev_num == dev);
+    }
+}
+
+int main()
+{
+  int res = -1;
+  str v2(2);
+
+#pragma omp target map(from:res)
+  {
+    res = v1.add(v2);
+  }
+
+  assert (res == 7);
+  check_host();
+  check_devs();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
new file mode 100644
index 0000000..6dd4260
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C
@@ -0,0 +1,50 @@
+// { dg-do run }
+// { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+// { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+// { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "_GLOBAL__off_I_v1" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "__omp_target_static_init_and_destruction" 2 "gimple" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 2 "gimple" } }
+
+// { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+// { dg-final { scan-tree-dump-not "__omp_target_static_init_and_destruction" "optimized" } }
+// FIXME: should be '-not' not '-times' 1:
+// { dg-final { scan-tree-dump-times "void _GLOBAL__off_I_" 1 "optimized" } }
+// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 1 "optimized" } }
+
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_amdgcn } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_nvptx } } }
+// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_nvptx } } }
+
+
+#include <cassert>
+
+#pragma omp declare target
+
+template<typename T>
+struct str {
+  str(T x) : _x(x) { }
+  T add(str o) { return _x + o._x; }
+  T _x;
+};
+
+str<long> v1(5);
+
+#pragma omp end declare target
+
+int main()
+{
+  long res = -1;
+  str<long> v2(2);
+
+#pragma omp target map(from:res)
+  {
+    res = v1.add(v2);
+  }
+
+  assert (res == 7);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C
new file mode 100644
index 0000000..8d4aff2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C
@@ -0,0 +1,36 @@
+// { dg-do run }
+
+#include <cassert>
+
+#pragma omp declare target
+
+struct item {
+  item(item *p, int v) : prev(p), val(v) { }
+  int get() { return prev ? prev->get() * val : val; }
+  item *prev;
+  int val;
+};
+
+/* This case demonstrates why constructing on the host and then copying to
+   the target would be less desirable.  With on-target construction, "prev"
+   for each 'item' will be a device pointer, not a host pointer.  */
+item hubert1(nullptr, 3);
+item hubert2(&hubert1, 5);
+item hubert3(&hubert2, 7);
+item hubert4(&hubert3, 11);
+
+#pragma omp end declare target
+
+int main()
+{
+  int res = -1;
+
+#pragma omp target map(from:res)
+  {
+    res = hubert4.get ();
+  }
+
+  assert (res == 1155);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c
new file mode 100644
index 0000000..313d188
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c
@@ -0,0 +1,43 @@
+/* { dg-do run } */
+/* { dg-additional-options "-fno-builtin-omp_is_initial_device" } */
+
+/* Check whether 'omp_is_initial_device()' is NOT compile-time optimized. */
+
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }  */
+/* { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }  */
+
+/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }  */
+
+/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "optimized" } }  */
+
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_amdgcn } } }  */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_nvptx } } }  */
+
+
+#include <omp.h>
+
+int
+main ()
+{
+  int is_initial, dev_num, initial;
+  initial = omp_get_initial_device();
+  for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+    {
+      is_initial = dev_num = 99;
+      #pragma omp target map(from: is_initial, dev_num) device(dev)
+        {
+          is_initial = omp_is_initial_device ();
+          dev_num = omp_get_device_num ();
+        }
+      if (dev == omp_initial_device || dev == initial)
+	{
+	  if (dev_num != initial || is_initial != 1)
+	    __builtin_abort ();
+	}
+      else
+	{
+	  if (dev_num != dev || is_initial != 0)
+	    __builtin_abort ();
+	}
+    }
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c
new file mode 100644
index 0000000..423727c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c
@@ -0,0 +1,42 @@
+/* { dg-do run } */
+
+/* Check whether 'omp_is_initial_device()' is properly compile-time optimized. */
+
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }  */
+/* { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }  */
+
+/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }  */
+
+/* { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }  */
+
+/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } }  */
+/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } }  */
+
+
+#include <omp.h>
+
+int
+main ()
+{
+  int is_initial, dev_num, initial;
+  initial = omp_get_initial_device();
+  for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+    {
+      is_initial = dev_num = 99;
+      #pragma omp target map(from: is_initial, dev_num) device(dev)
+        {
+          is_initial = omp_is_initial_device ();
+          dev_num = omp_get_device_num ();
+        }
+      if (dev == omp_initial_device || dev == initial)
+	{
+	  if (dev_num != initial || is_initial != 1)
+	    __builtin_abort ();
+	}
+      else
+	{
+	  if (dev_num != dev || is_initial != 0)
+	    __builtin_abort ();
+	}
+    }
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90 b/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90
new file mode 100644
index 0000000..e06ced2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90
@@ -0,0 +1,37 @@
+! { dg-additional-options "-fno-builtin-omp_is_initial_device" }
+
+! Check whether 'omp_is_initial_device()' is NOT compile-time optimized. */
+
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "optimized" } }
+
+! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_amdgcn } } }
+! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_nvptx } } }
+
+
+program main
+  use omp_lib
+  implicit none (type, external)
+  integer :: dev_num, initial, dev
+  logical :: is_initial
+
+  initial = omp_get_initial_device()
+  do dev = omp_initial_device, omp_get_num_devices()
+      dev_num = 99
+      !$omp target map(from: is_initial, dev_num) device(dev)
+        is_initial = omp_is_initial_device ()
+        dev_num = omp_get_device_num ()
+      !$omp end target
+      if (dev == omp_initial_device .or. dev == initial) then
+        if (dev_num /= initial .or. .not. is_initial) &
+          stop 1
+      else
+        if (dev_num /= dev .or. is_initial) &
+          stop 2
+      end if
+  end do
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f
new file mode 100644
index 0000000..fec4a3f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f
@@ -0,0 +1,35 @@
+! Check whether 'omp_is_initial_device()' is properly compile-time optimized. */
+
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+
+! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } }
+! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } }
+
+
+      program main
+      implicit none (type, external)
+      include "omp_lib.h" 
+      integer :: dev_num, initial, dev
+      logical :: is_initial
+
+      initial = omp_get_initial_device()
+      do dev = omp_initial_device, omp_get_num_devices()
+        dev_num = 99
+!$omp target map(from: is_initial, dev_num) device(dev)
+          is_initial = omp_is_initial_device ()
+          dev_num = omp_get_device_num ()
+!$omp end target
+        if (dev == omp_initial_device .or. dev == initial) then
+          if (dev_num /= initial .or. .not. is_initial)                         &
+     &      stop 1
+        else
+          if (dev_num /= dev .or. is_initial)                                   &
+     &      stop 2
+       end if
+      end do
+      end
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90 b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90
new file mode 100644
index 0000000..f8a645f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90
@@ -0,0 +1,35 @@
+! Check whether 'omp_is_initial_device()' is properly compile-time optimized. */
+
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" }
+! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } }
+
+! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } }
+
+! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } }
+! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } }
+
+
+program main
+  use omp_lib
+  implicit none (type, external)
+  integer :: dev_num, initial, dev
+  logical :: is_initial
+
+  initial = omp_get_initial_device()
+  do dev = omp_initial_device, omp_get_num_devices()
+      dev_num = 99
+      !$omp target map(from: is_initial, dev_num) device(dev)
+        is_initial = omp_is_initial_device ()
+        dev_num = omp_get_device_num ()
+      !$omp end target
+      if (dev == omp_initial_device .or. dev == initial) then
+        if (dev_num /= initial .or. .not. is_initial) &
+          stop 1
+      else
+        if (dev_num /= dev .or. is_initial) &
+          stop 2
+      end if
+  end do
+end