diff mbox series

[committed,OG10] Permit calls to builtins and intrinsics in kernels loops

Message ID 367f63c3-0164-79b9-5825-10bea1881d64@codesourcery.com
State New
Headers show
Series [committed,OG10] Permit calls to builtins and intrinsics in kernels loops | expand

Commit Message

Sandra Loosemore Aug. 23, 2020, 2:14 a.m. UTC 
This is another small tweak to the kernels loop annotator.  Like the 
subject line says, it allows the annotator to consider loops that have 
calls to builtins (C/C++) or intrinsics (Fortran).  I've committed this 
to the OG10 branch since there is no one to review these patches right 
now, and it will be mashed into the rest of the kernels loop annotation 
code for review before submission for mainline.

-Sandra
diff mbox series

Patch

commit 1c9af55d7ff76e2e6b633af33e6e6991a0ba4c48
Author: Sandra Loosemore <sandra@codesourcery.com>
Date:   Sat Aug 22 18:23:26 2020 -0700

    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 51a3b9e..2c48153 100644
--- 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 24f2448..34523cee 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2850,8 +2850,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;
@@ -2865,8 +2866,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 13cc5a6..df5448d 100644
--- 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-21  Tobias Burnus  <tobias@codesourcery.com>
 
 	Backport from mainline
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index d2e7b73..b297f4b 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -7094,9 +7094,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 c5f2813..4adda2a 100644
--- 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>
 
 	Backport from mainline
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 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
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" } }