From: Sandra Loosemore Date: Sun, 23 Aug 2020 01:23:26 +0000 (-0700) Subject: Permit calls to builtins and intrinsics in kernels loops. X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=b813e919ddc1eecb2e2c28002e775c2aecddd708;p=thirdparty%2Fgcc.git 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 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 3961227eb2ac..4bf22985b04d 100644 --- a/gcc/c-family/ChangeLog.omp +++ b/gcc/c-family/ChangeLog.omp @@ -1,3 +1,11 @@ +2020-08-22 Sandra Loosemore + + 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 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 7d2b69f36324..f7c5d845c4e3 100644 --- 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 e2a8b51f7ecb..f1536a5ff79f 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,11 @@ +2020-08-22 Sandra Loosemore + + 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 Annotate inner loops in "acc kernels loop" directives (Fortran). diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 7b0a5937ed59..6b18f77fc29b 100644 --- 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 cb274060614e..7ef9c3d6c9b7 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,11 @@ +2020-08-22 Sandra Loosemore + + 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 * 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 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" } }