[gcc/devel/omp/gcc-12] libgomp: Add offload_device_gcn check, add requires-4a.c test

[gcc/devel/omp/gcc-12] libgomp: Add offload_device_gcn check, add requires-4a.c test

[Tobias Burnus]

https://gcc.gnu.org/g:ffc6267a829068f49fd4346a0495d3f3fc51c89e

commit ffc6267a829068f49fd4346a0495d3f3fc51c89e
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Thu Oct 20 13:07:37 2022 +0200

    libgomp: Add offload_device_gcn check, add requires-4a.c test
    
    Duplicate libgomp.c-c++-common/requires-4.c (as ...-4a.c) but
    with using a heap-allocated instead of static memory for a variable.
    
    This change and the added offload_device_gcn check prepare for
    pseudo-USM, where the device hardware cannot access all host
    memory but only managed and pinned memory; for those, requires-4.c
    will fail and the new check permits to add
      target { ! { offload_device_nvptx || offload_device_gcn } }
    to requires-4.c; however, it has not been added yet as pseuo-USM
    support is not yet on mainline. (Review is pending for the USM
    patches.)
    
    include/ChangeLog:
    
            * gomp-constants.h (GOMP_DEVICE_HSA): Comment out unused define.
    
    libgomp/ChangeLog:
    
            * testsuite/lib/libgomp.exp (check_effective_target_offload_device_gcn):
            New.
            * testsuite/libgomp.c-c++-common/on_device_arch.h (device_arch_gcn,
            on_device_arch_gcn): New.
            * testsuite/libgomp.c-c++-common/requires-4a.c: New test; copied from
            requires-4.c but using heap-allocated memory.
    
    (cherry picked from commit 12d9f5afbd2660862045acd41cb65a77e35bea4d)

Diff:
---
 include/ChangeLog.omp                              |  7 ++++
 include/gomp-constants.h                           |  2 +-
 libgomp/ChangeLog.omp                              | 12 +++++++
 libgomp/testsuite/lib/libgomp.exp                  | 12 +++++++
 .../libgomp.c-c++-common/on_device_arch.h          | 13 ++++++++
 .../testsuite/libgomp.c-c++-common/requires-4a.c   | 39 ++++++++++++++++++++++
 6 files changed, 84 insertions(+), 1 deletion(-)

diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp
index 141ac55cf7a..4a72c8b9af5 100644
--- a/include/ChangeLog.omp
+++ b/include/ChangeLog.omp
@@ -1,3 +1,10 @@
+2022-10-20  Tobias Burnus  <tobias@codesourcery.com>
+
+	Backport from mainline:
+	2022-10-20  Tobias Burnus  <tobias@codesourcery.com>
+
+	* gomp-constants.h (GOMP_DEVICE_HSA): Comment out unused define.
+
 2022-08-31  Tobias Burnus  <tobias@codesourcery.com>
 
 	Revert:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 918414ec218..dd753a82106 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -257,7 +257,7 @@ enum gomp_map_kind
 #define GOMP_DEVICE_NOT_HOST		4
 #define GOMP_DEVICE_NVIDIA_PTX		5
 #define GOMP_DEVICE_INTEL_MIC		6
-#define GOMP_DEVICE_HSA			7
+/* #define GOMP_DEVICE_HSA		7 removed.  */
 #define GOMP_DEVICE_GCN			8
 
 /* We have a compatibility issue.  OpenMP 5.2 introduced
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 570aabf82c4..cb884aa0c83 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,15 @@
+2022-10-20  Tobias Burnus  <tobias@codesourcery.com>
+
+	Backport from mainline:
+	2022-10-20  Tobias Burnus  <tobias@codesourcery.com>
+
+	* testsuite/lib/libgomp.exp (check_effective_target_offload_device_gcn):
+	New.
+	* testsuite/libgomp.c-c++-common/on_device_arch.h (device_arch_gcn,
+	on_device_arch_gcn): New.
+	* testsuite/libgomp.c-c++-common/requires-4a.c: New test; copied from
+	requires-4.c but using heap-allocated memory.
+
 2022-10-17  Tobias Burnus  <tobias@codesourcery.com>
 
 	Backport from mainline:
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index c98c1edf57b..ffdcfd85b72 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -414,6 +414,18 @@ proc check_effective_target_offload_device_nvptx { } {
     } ]
 }
 
+# Return 1 if using a GCN offload device.
+proc check_effective_target_offload_device_gcn { } {
+    return [check_runtime_nocache offload_device_gcn {
+      #include <omp.h>
+      #include "testsuite/libgomp.c-c++-common/on_device_arch.h"
+      int main ()
+	{
+	  return !on_device_arch_gcn ();
+	}
+    } ]
+}
+
 # Return 1 if at least one Nvidia GPU is accessible.
 
 proc check_effective_target_openacc_nvidia_accel_present { } {
diff --git a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
index f92743b04d7..6f66dbd784c 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
+++ b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
@@ -7,6 +7,12 @@ device_arch_nvptx (void)
   return GOMP_DEVICE_NVIDIA_PTX;
 }
 
+/* static */ int
+device_arch_gcn (void)
+{
+  return GOMP_DEVICE_GCN;
+}
+
 /* static */ int
 device_arch_intel_mic (void)
 {
@@ -14,6 +20,7 @@ device_arch_intel_mic (void)
 }
 
 #pragma omp declare variant (device_arch_nvptx) match(construct={target},device={arch(nvptx)})
+#pragma omp declare variant (device_arch_gcn) match(construct={target},device={arch(gcn)})
 #pragma omp declare variant (device_arch_intel_mic) match(construct={target},device={arch(intel_mic)})
 /* static */ int
 device_arch (void)
@@ -37,6 +44,12 @@ on_device_arch_nvptx ()
   return on_device_arch (GOMP_DEVICE_NVIDIA_PTX);
 }
 
+int
+on_device_arch_gcn ()
+{
+  return on_device_arch (GOMP_DEVICE_GCN);
+}
+
 int
 on_device_arch_intel_mic ()
 {
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c
new file mode 100644
index 00000000000..4fb9783a97a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c
@@ -0,0 +1,39 @@
+/* { dg-additional-options "-flto" } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+/* { dg-additional-sources requires-4-aux.c } */
+
+/* Same as requires-4.c, but uses heap memory for 'a'.  */
+
+/* Check no diagnostic by device-compiler's or host compiler's lto1.
+   Other file uses: 'requires reverse_offload', but that's inactive as
+   there are no declare target directives, device constructs nor device routines  */
+
+/* Depending on offload device capabilities, it may print something like the
+   following (only) if GOMP_DEBUG=1:
+   "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled"
+   and in that case does host-fallback execution.
+
+   No offload devices support USM at present, so we may verify host-fallback
+   execution by presence of separate memory spaces.  */
+
+#pragma omp requires unified_address,unified_shared_memory
+
+int *a;
+extern void foo (void);
+
+int
+main (void)
+{
+  a = (int *) __builtin_calloc (sizeof (int), 10);
+  #pragma omp target map(to: a)
+  for (int i = 0; i < 10; i++)
+    a[i] = i;
+
+  for (int i = 0; i < 10; i++)
+    if (a[i] != i)
+      __builtin_abort ();
+
+  foo ();
+  __builtin_free (a);
+  return 0;
+}