diff mbox series

[09/40] Permit calls to builtins and intrinsics in kernels loops.

Message ID 20211215155447.19379-10-frederik@codesourcery.com
State New
Headers show
Series OpenACC "kernels" Improvements | expand

Commit Message

Frederik Harwath Dec. 15, 2021, 3:54 p.m. UTC 
From: Sandra Loosemore <sandra@codesourcery.com>

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.
---
 gcc/c-family/c-omp.c                          | 10 ++++---
 gcc/fortran/openmp.c                          |  9 ++++---
 .../goacc/kernels-loop-annotation-20.c        | 23 ++++++++++++++++
 .../goacc/kernels-loop-annotation-20.f95      | 26 +++++++++++++++++++
 4 files changed, 61 insertions(+), 7 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-loop-annotation-20.c
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-20.f95

--
2.33.0

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
diff mbox series

Patch

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 30757877eafe..e7c27f45e888 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3545,8 +3545,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;
@@ -3560,8 +3561,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/openmp.c b/gcc/fortran/openmp.c
index b0b68b494778..d5d996e378d7 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -9156,9 +9156,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/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
index 000000000000..5e3f02845713
--- /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
index 000000000000..5169a0a1676d
--- /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" } }