From patchwork Wed Dec 15 15:54:16 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 48951 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 09A36385741B for ; Wed, 15 Dec 2021 16:01:14 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 04A04385800C; Wed, 15 Dec 2021 15:55:36 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 04A04385800C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: WXfWwOSvGMJtzcbSW7y6guXw2qETziFk+MTqnfA9sE9odgbzt2oR+wAu7als30vuktA7XwM4CH J3VIMsmR/FIt/i2Gu9Cl0OLspeo2/irzLrZ8AqSPRER9IgIWlA/Zhs5B/CHcFP0cRWk2pba7A9 07Wr6O8EcTueClgblszCUnlLh4u7LB26VccUryMTHSU8v8KbWeBCJBjUN99QIJNZ3bXLPN+UvM 6yfMkFk5RE1FNC9C8MBmIuf8wosjqy/DTM5rUzNIaGLVrRBN9tM7XiKY2ebHKRKNKbRrxvYIoo 1wMl6JyykgsodUbG6d9kZ7he X-IronPort-AV: E=Sophos;i="5.88,207,1635235200"; d="scan'208";a="69738358" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 15 Dec 2021 07:55:33 -0800 IronPort-SDR: fTundpMeqxVVh1lMKpSA1O9M6PaZBgTqj2wVSLn0ua2oafHzjFKLfqF5TdTUg4IxdpM5vcW2nf rY9u0oK9Hk8DBAcCttXbjN+woHo1ngBDKRHn5M0GxtrliLFCJBbD0Mp5Cd2JVrqRqKjMb2SaXV mYKjVciERR6688IXeDLEzz+CWToRsqbvHqxAOAxrBXro4p1gcNnz9APODTZDOIJlhO+k5CsUwB dGvrb2YKt6R1sXgap429fSllx7N/0tWGI3n2wWpf9dW9G3c/US4/ill0ANCzXIxdRm1r3eyCez O9c= From: Frederik Harwath To: Subject: [PATCH 09/40] Permit calls to builtins and intrinsics in kernels loops. Date: Wed, 15 Dec 2021 16:54:16 +0100 Message-ID: <20211215155447.19379-10-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211215155447.19379-1-frederik@codesourcery.com> References: <20211215155447.19379-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: fortran@gcc.gnu.org, nathan@acm.org, Sandra Loosemore , tobias@codesourcery.com, thomas@codesourcery.com, joseph@codesourcery.com Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" From: Sandra Loosemore 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.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 --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" } }