OpenACC 'nohost' clause

Do not "compile a version of this procedure for the host".

	gcc/
	* tree-core.h (omp_clause_code): Add 'OMP_CLAUSE_NOHOST'.
	* tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
	Handle it.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* omp-general.c (oacc_verify_routine_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses)
	(gimplify_adjust_omp_clauses): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses)
	(convert_local_omp_clauses): Likewise.
	* omp-low.c (scan_sharing_clauses): Likewise.
	* omp-offload.c (execute_oacc_device_lower): Update.
	gcc/c-family/
	* c-pragma.h (pragma_omp_clause): Add 'PRAGMA_OACC_CLAUSE_NOHOST'.
	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Handle 'nohost'.
	(c_parser_oacc_all_clauses): Handle 'PRAGMA_OACC_CLAUSE_NOHOST'.
	(OACC_ROUTINE_CLAUSE_MASK): Add 'PRAGMA_OACC_CLAUSE_NOHOST'.
	* c-typeck.c (c_finish_omp_clauses): Handle 'OMP_CLAUSE_NOHOST'.
	gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Handle 'nohost'.
	(cp_parser_oacc_all_clauses): Handle 'PRAGMA_OACC_CLAUSE_NOHOST'.
	(OACC_ROUTINE_CLAUSE_MASK): Add 'PRAGMA_OACC_CLAUSE_NOHOST'.
	* pt.c (tsubst_omp_clauses): Handle 'OMP_CLAUSE_NOHOST'.
	* semantics.c (finish_omp_clauses): Likewise.
	gcc/fortran/
	* dump-parse-tree.c (show_attr): Update.
	* gfortran.h (symbol_attribute): Add 'oacc_routine_nohost' member.
	(gfc_omp_clauses): Add 'nohost' member.
	* module.c (ab_attribute): Add 'AB_OACC_ROUTINE_NOHOST'.
	(attr_bits, mio_symbol_attribute): Update.
	* openmp.c (omp_mask2): Add 'OMP_CLAUSE_NOHOST'.
	(gfc_match_omp_clauses): Handle 'OMP_CLAUSE_NOHOST'.
	(OACC_ROUTINE_CLAUSES): Add 'OMP_CLAUSE_NOHOST'.
	(gfc_match_oacc_routine): Update.
	* trans-decl.c (add_attributes_to_decl): Update.
	* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
	gcc/testsuite/
	* c-c++-common/goacc/classify-routine-nohost.c: New file.
	* c-c++-common/goacc/classify-routine.c: Update.
	* c-c++-common/goacc/routine-2.c: Likewise.
	* c-c++-common/goacc/routine-nohost-1.c: New file.
	* c-c++-common/goacc/routine-nohost-2.c: Likewise.
	* g++.dg/goacc/template.C: Update.
	* gfortran.dg/goacc/classify-routine-nohost.f95: New file.
	* gfortran.dg/goacc/classify-routine.f95: Update.
	* gfortran.dg/goacc/pure-elemental-procedures-2.f90: Likewise.
	* gfortran.dg/goacc/routine-6.f90: Likewise.
	* gfortran.dg/goacc/routine-intrinsic-2.f: Likewise.
	* gfortran.dg/goacc/routine-module-1.f90: Likewise.
	* gfortran.dg/goacc/routine-module-2.f90: Likewise.
	* gfortran.dg/goacc/routine-module-3.f90: Likewise.
	* gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
	* gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise.
	* gfortran.dg/goacc/routine-multiple-directives-2.f90: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c:
	Likewise.
	* testsuite/libgomp.oacc-fortran/routine-nohost-1.f90: Likewise.

Co-Authored-By: Joseph Myers <joseph@codesourcery.com>
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

4 files changed, 183 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
new file mode 100644
index 0000000..dc92727
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
@@ -0,0 +1,63 @@
+/* Test 'nohost' clause via 'acc_on_device'.
+
+   With optimizations disabled, we currently don't expect that 'acc_on_device' "evaluates at compile time to a constant".
+   { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } }
+*/
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+/* { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'.  */
+
+#include <assert.h>
+#include <openacc.h>
+
+#pragma acc routine
+static int fact(int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+  else
+    return n * fact(n - 1);
+}
+
+#pragma acc routine nohost
+static int fact_nohost(int n)
+{
+  /* Make sure this fails host compilation.  */
+#if defined ACC_DEVICE_TYPE_host
+  asm ("IT'S A TRAP");
+#elif defined ACC_DEVICE_TYPE_nvidia
+  asm ("{\n\t  .reg .u32 %tid_x;\n\t  mov.u32 %tid_x, %tid.x;\n\t}");
+#elif defined ACC_DEVICE_TYPE_radeon
+  asm ("s_nop 0");
+#else
+# error Not ported to this ACC_DEVICE_TYPE
+#endif
+
+  return fact(n);
+}
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target c } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'int fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && offloading_enabled } } } }
+   TODO See PR101551 for 'offloading_enabled' differences.  */
+
+int main()
+{
+#define N 10
+  int x[N];
+
+#pragma acc parallel loop copyout(x)
+  for (int i = 0; i < N; ++i)
+    /*TODO PR82391: '(int) acc_device_*' cast to avoid the C++ 'acc_on_device' wrapper.  */
+    x[i] = acc_on_device((int) acc_device_not_host) ? fact_nohost(i) : 0;
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (acc_get_device_type() == acc_device_host)
+	assert(x[i] == 0);
+      else
+	assert(x[i] == fact(i));
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c
new file mode 100644
index 0000000..4d081f2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c
@@ -0,0 +1,39 @@
+/* Test 'nohost' clause via 'weak'.
+
+   { dg-require-effective-target weak_undefined }
+
+   When the OpenACC 'routine' with 'nohost' clauses gets discarded, the weak symbol then resolves to 'NULL'.
+*/
+
+/* { dg-additional-sources routine-nohost-2_2.c } */
+
+/* { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'.  */
+
+#include <assert.h>
+#include <openacc.h>
+
+#pragma acc routine //nohost
+__attribute__((weak))
+extern int f1(int);
+
+int main()
+{
+  int x = -10;
+
+#pragma acc serial copy(x)
+  /* { dg-warning {using vector_length \(32\), ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } */
+  {
+    if (f1)
+      x = f1(x);
+    else
+      x = 0;
+
+  }
+
+  if (acc_get_device_type() == acc_device_host)
+    assert(x == 0);
+  else
+    assert(x == -20);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c
new file mode 100644
index 0000000..6029545
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-2_2.c
@@ -0,0 +1,18 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma acc routine nohost
+int f1(int x)
+{
+  /* Make sure this fails host compilation.  */
+#if defined ACC_DEVICE_TYPE_host
+  asm ("IT'S A TRAP");
+#elif defined ACC_DEVICE_TYPE_nvidia
+  asm ("{\n\t  .reg .u32 %tid_x;\n\t  mov.u32 %tid_x, %tid.x;\n\t}");
+#elif defined ACC_DEVICE_TYPE_radeon
+  asm ("s_nop 0");
+#else
+# error Not ported to this ACC_DEVICE_TYPE
+#endif
+
+  return 2 * x;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90
new file mode 100644
index 0000000..cd5bddc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90
@@ -0,0 +1,63 @@
+! Test 'nohost' clause via 'acc_on_device'.
+
+! { dg-do run }
+
+! With optimizations disabled, we currently don't expect that 'acc_on_device' "evaluates at compile time to a constant".
+! { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } }
+
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 10
+  integer :: a(n), i
+  integer, external :: fact_nohost
+  !$acc routine (fact_nohost)
+  integer, external :: fact
+
+  !$acc parallel loop
+  do i = 1, n
+     if (acc_on_device(acc_device_not_host)) then
+        a(i) = fact_nohost(i)
+     else
+        a(i) = 0
+     end if
+  end do
+  !$acc end parallel loop
+
+  do i = 1, n
+     if (acc_get_device_type() .eq. acc_device_host) then
+        if (a(i) .ne. 0) stop 10 + i
+     else
+        if (a(i) .ne. fact(i)) stop 20 + i
+     end if
+  end do
+end program main
+
+recursive function fact(x) result(res)
+  implicit none
+  !$acc routine (fact)
+  integer, intent(in) :: x
+  integer :: res
+
+  if (x < 1) then
+     res = 1
+  else
+     res = x * fact(x - 1)
+  end if
+end function fact
+
+function fact_nohost(x) result(res)
+  use openacc
+  implicit none
+  !$acc routine (fact_nohost) nohost
+  integer, intent(in) :: x
+  integer :: res
+  integer, external :: fact
+
+  res = fact(x)
+end function fact_nohost
+! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost_' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
+!TODO See PR101551 for 'offloading_enabled' differences.