From 62312c778aaf7bc4b33d645188ec8a906fb5711b Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 12 Jul 2022 09:59:10 +0200
Subject: [PATCH] Add 'libgomp.c++/pr101544-1{,-O0}.C',
 'libgomp.oacc-c++/pr101544-1.C' [PR101544]

	PR target/101544
	libgomp/
	* testsuite/libgomp.c++/pr101544-1.C: New.
	* testsuite/libgomp.c++/pr101544-1-O0.C: Likewise.
	* testsuite/libgomp.oacc-c++/pr101544-1.C: Likewise.
---
 libgomp/testsuite/libgomp.c++/pr101544-1-O0.C |  4 +
 libgomp/testsuite/libgomp.c++/pr101544-1.C    | 83 +++++++++++++++++++
 .../testsuite/libgomp.oacc-c++/pr101544-1.C   |  7 ++
 3 files changed, 94 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.c++/pr101544-1-O0.C
 create mode 100644 libgomp/testsuite/libgomp.c++/pr101544-1.C
 create mode 100644 libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C

diff --git a/libgomp/testsuite/libgomp.c++/pr101544-1-O0.C b/libgomp/testsuite/libgomp.c++/pr101544-1-O0.C
new file mode 100644
index 00000000000..19a8cad2328
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr101544-1-O0.C
@@ -0,0 +1,4 @@
+// { dg-additional-options -foffload-options=-lstdc++ }
+// { dg-additional-options -O0 }
+
+#include "pr101544-1.C"
diff --git a/libgomp/testsuite/libgomp.c++/pr101544-1.C b/libgomp/testsuite/libgomp.c++/pr101544-1.C
new file mode 100644
index 00000000000..8ea70a242bc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr101544-1.C
@@ -0,0 +1,83 @@
+// { dg-additional-options -foffload-options=-lstdc++ }
+// See also '../libgomp.oacc-c++/pr101544-1.C'.
+#ifndef ALWAYS_INLINE
+# define ALWAYS_INLINE
+#endif
+
+//===--- declare_target_base_class.cpp --------------------------------------===//
+// 
+// OpenMP API Version 4.5 Nov 2015
+//
+// This test was suggested by members of NERSC. This test defines a declare
+// target region which includes only a base class and a 'concrete' device
+// pointer. 
+// 
+// Test suggestion comes from Chris Daily and Rahulkumar Gayatri from NERSC
+////===----------------------------------------------------------------------===//
+
+#include <new>
+#include <vector>
+#include <iostream>
+
+#pragma omp declare target
+//#pragma acc routine //TODO error: '#pragma acc routine' not immediately followed by function declaration or definition
+class S {
+public:
+  //#pragma acc routine //TODO error: '#pragma acc routine' must be at file scope
+  ALWAYS_INLINE
+  S() : _devPtr(nullptr) {}
+  //#pragma acc routine //TODO error: '#pragma acc routine' must be at file scope
+  ALWAYS_INLINE
+  double sag(double x, double y) {
+    return x + y;
+  }
+  S* cloneToDevice() {
+    S* ptr;
+#pragma omp target map(ptr)
+#pragma acc serial copy(ptr)
+    {
+      ptr = new S();
+    }
+    _devPtr = ptr;
+    return ptr;
+  }
+private:
+  S* _devPtr;
+};
+//#pragma acc routine (S) //TODO error: 'class S' does not refer to a function
+//#pragma acc routine (S::S) //TODO error: '#pragma acc routine' names a set of overloads
+//#pragma acc routine (S::sag) //TODO error: '#pragma acc routine' names a set of overloads
+#pragma omp end declare target
+
+int main() {
+  int errors = 0;
+  
+  S s;
+  S* devPtr = s.cloneToDevice();
+
+  std::vector<double> in(10, 0.0);
+  for(int i = 0; i < 10; i++) {
+    in[i] = i;
+  }
+
+  std::vector<double> out(10, 0.0);
+
+  double* inptr = in.data();
+  double* outptr = out.data();
+
+#pragma omp target teams distribute parallel for map(inptr[:10], outptr[:10]) is_device_ptr(devPtr)
+#pragma acc parallel loop copy(inptr[:10], outptr[:10]) deviceptr(devPtr)
+  for(int i = 0; i < 10; i++) {
+    outptr[i] = devPtr->sag(inptr[i], inptr[i]);
+  }
+
+  for(int i = 0; i < 10; i++) {
+    if (out[i] != i * 2)
+      {
+	++errors;
+	std::cerr << "ERROR: " << "i = " << i << ": " << out[i] << " != " << (i * 2) << "\n";
+      }
+  }
+
+  return errors ? 1 : 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C b/libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C
new file mode 100644
index 00000000000..004cd49353b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C
@@ -0,0 +1,7 @@
+// { dg-additional-options -foffload-options=-lstdc++ }
+// { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+// But actually, as none of the '#pragma acc routine' syntax is accepted, force inlining:
+#define ALWAYS_INLINE __attribute__((always_inline))
+
+#include "../libgomp.c++/pr101544-1.C"
+//TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} }