From patchwork Sun Aug 23 02:14:25 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 1349791 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4BYzNL1gg7z9sTS for ; Sun, 23 Aug 2020 12:14:41 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9E40E3857C75; Sun, 23 Aug 2020 02:14:37 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id E15783858D35 for ; Sun, 23 Aug 2020 02:14:33 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org E15783858D35 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Sandra_Loosemore@mentor.com IronPort-SDR: H+43YZVzJUwyUtfNN5l4XjXwO6ycMLgH5/ZdU1SoFQszJ/neaDInj0x0clSquC+Hsm6h7HCAvN QxKcU/Pq03cSBgApXxAhRIGBH5U3jJ6Khic6zEwVTeeT3rLWdNFRQ7DTRaxK1df7RCkFfP+ntD jtOOVoCl/K5gygGTzDw6DhzKepU2xpCIHpqSY2ECJIo0suVxqtOI8MoVoaKUTyQhneXle9l94e 8MfZWIWMgRMr0bUuyxdthbuAfCjU6U3Bau/hmgyBfrFARMNzE+T9kKQTWsPHsZJuM/a2bZdSLH EVw= X-IronPort-AV: E=Sophos;i="5.76,342,1592899200"; d="scan'208";a="54349344" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 22 Aug 2020 18:14:32 -0800 IronPort-SDR: HXQR44ZjpxZiBUZOsv0I4kMqhYD4qCqzpqyGTYGBzYGcz06t70zh6ll5/kzMwhwNlX5sjchuaM 1tAqNzPVFjQByIUnjeFVYXHKMvl6xZIcn0s8maokIrIvZyhLy4Y1NyvNcP97Kf8rFIJmQBn52r vHrgzbGBdHJbmn5yMOtjLWACBO8btKU1FNFXqc4PP4UXh+9Bcdf3fdwTvXkeOXmySglyZrr38I PMjBuvnxZhuyBONMDcykzF8jQTZVeS5j19PG+CYopOZfmlQm+BcHw2L+W5NZWiZ3pzV9kiqx3l jKo= To: "gcc-patches@gcc.gnu.org" From: Sandra Loosemore Subject: [committed] [OG10] Permit calls to builtins and intrinsics in kernels loops Message-ID: <367f63c3-0164-79b9-5825-10bea1881d64@codesourcery.com> Date: Sat, 22 Aug 2020 20:14:25 -0600 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.8.0 MIME-Version: 1.0 Content-Language: en-US X-ClientProxiedBy: svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) To svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) X-Spam-Status: No, score=-9.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.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: , Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" 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 commit 1c9af55d7ff76e2e6b633af33e6e6991a0ba4c48 Author: Sandra Loosemore 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 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 + + 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 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 + + 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 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 + + 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 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" } }