Permit calls to builtins and intrinsics in kernels loops.

This tweak to the OpenACC kernels loop annotation relaxes the
restrictions on function calls in the loop body.  Normally calls to
functions not explicitly marked with a parallelism attribute are not
permitted, but C/C++ builtins and Fortran intrinsics have known
semantics so we can generally permit those without restriction.  If
any turn out to be problematical, we can add on here to recognize
them, or in the processing of the "auto" annotations.

2020-08-22  Sandra Loosemore  <sandra@codesourcery.com>

	gcc/c-family/
	* c-omp.c (annotate_loops_in_kernels_regions): Test for
	calls to builtins.

	gcc/fortran/
	* openmp.c (check_expr_for_invalid_calls): Check for intrinsic
	functions.

	gcc/testsuite/
	* c-c++-common/goacc/kernels-loop-annotation-20.c: New.
	* gfortran.dg/goacc/kernels-loop-annotation-20.f95: New.

diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 90aa0e6a17d065fe5e274cf67be3804bf1286b94..4c99f6ea2285241bd07626768e86c80356e2071f 100644 (file)
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,3 +1,11 @@
+2020-08-22  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Allow annotation of loops containing calls to builtins in
+       kernels regions.
+
+       * c-omp.c (annotate_loops_in_kernels_regions): Test for
+       calls to builtins.
+
 2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
 
        Annotate inner loops in "acc kernels loop" directives (C/C++).

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 71875efff557524e49c36839a60077e0496e1530..1532462d263895d795b1717001518ebc37af6566 100644 (file)
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3317,8 +3317,9 @@ annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees,
       break;
 
     case CALL_EXPR:
-      /* Direct function calls to functions marked as OpenACC routines are
-        allowed.  Reject indirect calls or calls to non-routines.  */
+      /* Direct function calls to builtins and functions marked as
+        OpenACC routines are allowed.  Reject indirect calls or calls
+        to non-routines.  */
       if (info->state >= as_in_kernels_loop)
        {
          tree fn = CALL_EXPR_FN (node), fn_decl = NULL_TREE;
@@ -3332,8 +3333,9 @@ annotate_loops_in_kernels_regions (tree *nodeptr, int *walk_subtrees,
            }
          if (fn_decl == NULL_TREE)
            do_not_annotate_loop_nest (info, as_invalid_call, node);
-         else if (!lookup_attribute ("oacc function",
-                                     DECL_ATTRIBUTES (fn_decl)))
+         else if (!fndecl_built_in_p (fn_decl, BUILT_IN_NORMAL)
+                  && !lookup_attribute ("oacc function",
+                                        DECL_ATTRIBUTES (fn_decl)))
            do_not_annotate_loop_nest (info, as_invalid_call, node);
        }
       break;

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index a56ca22aefc73483624e2d821031ecc145f8ea60..48277068b6bc234462eed341871f2fa15568618e 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,11 @@
+2020-08-22  Sandra Loosemore  <sandra@codesourcery.com>
+
+       Permit calls to Fortran intrinsics when annotating loops in
+       kernels regions.
+
+       * openmp.c (check_expr_for_invalid_calls): Check for intrinsic
+       functions.
+
 2020-08-19  Sandra Loosemore  <sandra@codesourcery.com>
 
        Annotate inner loops in "acc kernels loop" directives (Fortran).

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 01626e50c0e6c5f407f0daf3f34dd0ec0b6d7def..dca2a782d7398fd4bb2f02678cc2386dbcc7ac23 100644 (file)
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -7525,9 +7525,12 @@ check_expr_for_invalid_calls (gfc_expr **exprp, int *walk_subtrees,
   switch (expr->expr_type)
     {
     case EXPR_FUNCTION:
-      if (expr->value.function.esym
-         && (expr->value.function.esym->attr.oacc_routine_lop
-             != OACC_ROUTINE_LOP_NONE))
+      /* Permit calls to Fortran intrinsic functions and to routines
+        with an explicitly declared parallelism level.  */
+      if (expr->value.function.isym
+         || (expr->value.function.esym
+             && (expr->value.function.esym->attr.oacc_routine_lop
+                 != OACC_ROUTINE_LOP_NONE)))
        return 0;
       /* Else fall through.  */
 

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index d5c60a404fa0ddcd622caaea43f254c676f8f161..7672b8d76f8c869d043ecf55393fc744ea31ae0c 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,11 @@
+2020-08-22   Sandra Loosemore  <sandra@codesourcery.com>
+
+       Test cases for allowing calls to C/C++ builtins and Fortran
+       intrinsics in loops in kernels regions.
+
+       * c-c++-common/goacc/kernels-loop-annotation-20.c: New.
+       * gfortran.dg/goacc/kernels-loop-annotation-20.f95: New.
+
 2020-08-21  Tobias Burnus  <tobias@codesourcery.com>
 
        * gfortran.dg/gomp/pr67500.f90: Change dg-warning to

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-20.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-20.c
new file mode 100644 (file)
index 0000000..5e3f028
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-20.c
@@ -0,0 +1,23 @@
+/* { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-Wopenacc-kernels-annotate-loops" } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-do compile } */
+
+/* Test that calls to built-in functions don't inhibit kernels loop
+   annotation.  */
+
+void foo (int n, int *input, int *out1, int *out2)
+{
+#pragma acc kernels
+  {
+    int i;
+
+    for (i = 0; i < n; i++)
+      {
+       out1[i] = __builtin_clz (input[i]);
+       out2[i] = __builtin_popcount (input[i]);
+      }
+  }
+}
+
+/* { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } } */

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95
new file mode 100644 (file)
index 0000000..5169a0a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95
@@ -0,0 +1,26 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with calls to intrinsics in the body can be annotated.
+
+subroutine f (n, input, out1, out2)
+  implicit none
+  integer :: n
+  integer, intent (in), dimension (n) :: input
+  integer, intent (out), dimension (n) :: out1, out2
+
+  integer :: i
+
+!$acc kernels
+
+  do i = 1, n
+      out1(i) = min (i, input(i))
+      out2(i) = not (input(i))
+  end do
+!$acc end kernels
+
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop auto" 1 "original" } }