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.cc (annotate_loops_in_kernels_regions): Test for
	calls to builtins.

	gcc/fortran/
	* openmp.cc (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 3961227eb2acc3a81950254ad9f3dbffb317c50c..4bf22985b04db76fa1cb222d9d3d208a051d1b17 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.cc (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.cc b/gcc/c-family/c-omp.cc
index 7d2b69f36324c733c5fcf087709b0e8f5a636f81..f7c5d845c4e393360507b797ee6bd020a89fd77f 100644 (file)
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3540,8 +3540,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;
@@ -3555,8 +3556,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 e2a8b51f7ecb35aaf55cfa3ebf702a61f4c208e5..f1536a5ff79f9c25b29ae9746569a637d24e5226 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.cc (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.cc b/gcc/fortran/openmp.cc
index 7b0a5937ed5961b9a9a008a6429a3b9d68041843..6b18f77fc29b295025f48271430b4db48949953c 100644 (file)
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -9530,9 +9530,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 cb274060614ea5eea63228fffb99e644c6375dd8..7ef9c3d6c9b777df88bd90790c0efabc080dde3d 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" } }