From patchwork Fri Nov 13 22:11:59 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1400160 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 4CXt4K6dnDz9sTc for ; Sat, 14 Nov 2020 09:12:17 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D154D3951015; Fri, 13 Nov 2020 22:12:15 +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 B90923857025 for ; Fri, 13 Nov 2020 22:12:10 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org B90923857025 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: lJkkCqPTn+7O7Gs7dqeYlVv58cgl+egU3SKFxX4a0BE1V60n3IEN+VS4kxyV5Fd7kAapvgnmxQ y640J4Uh5lSDnsKPmIHLG6eEOTpNtVEuwD88PijgaM3vYUNk4SLhhqwGJtqkcfbgFxP5GmrGqw LumAQWyR2NDPRpqxJfyPf+MNqokRkmvARgghSKjjyBHGh0+r2FDWX5ThyjgpYgbKtff5Fq6Abz exVC2aIUfc/uBLXTFjnIh6gZn0CO+SFbItpHmup81AiO2nhPzs5gOLUGEQWU84BdNho/29caTj a8M= X-IronPort-AV: E=Sophos;i="5.77,476,1596528000"; d="scan'208,223";a="55158141" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 13 Nov 2020 14:12:09 -0800 IronPort-SDR: 3A0ycgucSShZmcLcDH96u2COCwutSCvZM0nRyQZHU5I9+oARaLtbw/MeRW5qVhP/w4IVKN7tb+ bam2VROtFIimsAo/VzX7eyXh1baDuq5TDdWLCqzDECbNlt7D0YqfF6MKlQxTUybhgzsW9qE5zk oPzdmFkNhAAkwXRSlne8lyh5wSmIKC/Z1Qk+OjDZlCsuaetUeLyTuFXZKYUA/qAILixJ4PpNxH 4F4Ka7dBiiy5CUQN+oi6YpZNiEGxy+RgAnqYY4xHOIh4MlcEv2/Rd+vgpC1lo9RQxL+L7iBkmQ 9SM= From: Thomas Schwinge To: Subject: Attach an attribute to all outlined OpenACC compute regions User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Fri, 13 Nov 2020 23:11:59 +0100 Message-ID: <87eekwdhzk.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, 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" Hi! I've pushed "Attach an attribute to all outlined OpenACC compute regions" to master branch in commit 703e4f86496214e4915db898397fcd0ae1d955e0, and backported to releases/gcc-10 branch in commit 40bf92be5b621318a43347236508696cc387f3a6, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter From 40bf92be5b621318a43347236508696cc387f3a6 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 28 Oct 2020 11:43:49 +0100 Subject: [PATCH] Attach an attribute to all outlined OpenACC compute regions This allows for making some things more explicit, later on. gcc/ * omp-expand.c (expand_omp_target): Attach an attribute to all outlined OpenACC compute regions. * omp-offload.c (execute_oacc_device_lower): Adjust. gcc/testsuite/ * c-c++-common/goacc/classify-parallel.c: Adjust. * gfortran.dg/goacc/classify-parallel.f95: Likewise. * c-c++-common/goacc/classify-serial.c: New. * gfortran.dg/goacc/classify-serial.f95: Likewise. (cherry picked from commit 703e4f86496214e4915db898397fcd0ae1d955e0) --- gcc/omp-expand.c | 22 +++++--- gcc/omp-offload.c | 51 +++++++++++++------ .../c-c++-common/goacc/classify-parallel.c | 4 +- .../c-c++-common/goacc/classify-serial.c | 29 +++++++++++ .../gfortran.dg/goacc/classify-parallel.f95 | 4 +- .../gfortran.dg/goacc/classify-serial.f95 | 31 +++++++++++ 6 files changed, 114 insertions(+), 27 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-serial.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index da1f4c39d18..735e263c8f8 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7947,27 +7947,33 @@ expand_omp_target (struct omp_region *region) entry_bb = region->entry; exit_bb = region->exit; + if (target_kind == GF_OMP_TARGET_KIND_OACC_KERNELS) + mark_loops_in_oacc_kernels_region (region->entry, region->exit); + + /* Going on, all OpenACC compute constructs are mapped to + 'BUILT_IN_GOACC_PARALLEL', and get their compute regions outlined. + To distinguish between them, we attach attributes. */ switch (target_kind) { + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + DECL_ATTRIBUTES (child_fn) + = tree_cons (get_identifier ("oacc parallel"), + NULL_TREE, DECL_ATTRIBUTES (child_fn)); + break; case GF_OMP_TARGET_KIND_OACC_KERNELS: - mark_loops_in_oacc_kernels_region (region->entry, region->exit); - - /* Further down, all OpenACC compute constructs will be mapped to - BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there - is an "oacc kernels" attribute set for OpenACC kernels. */ DECL_ATTRIBUTES (child_fn) = tree_cons (get_identifier ("oacc kernels"), NULL_TREE, DECL_ATTRIBUTES (child_fn)); break; case GF_OMP_TARGET_KIND_OACC_SERIAL: - /* Further down, all OpenACC compute constructs will be mapped to - BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there - is an "oacc serial" attribute set for OpenACC serial. */ DECL_ATTRIBUTES (child_fn) = tree_cons (get_identifier ("oacc serial"), NULL_TREE, DECL_ATTRIBUTES (child_fn)); break; default: + /* Make sure we don't miss any. */ + gcc_checking_assert (!(is_gimple_omp_oacc (entry_stmt) + && is_gimple_omp_offloaded (entry_stmt))); break; } diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 6daf6226921..5ce3dffb562 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -1541,12 +1541,45 @@ execute_oacc_device_lower () flag_openacc_dims = (char *)&flag_openacc_dims; } + bool is_oacc_parallel + = (lookup_attribute ("oacc parallel", + DECL_ATTRIBUTES (current_function_decl)) != NULL); bool is_oacc_kernels = (lookup_attribute ("oacc kernels", DECL_ATTRIBUTES (current_function_decl)) != NULL); + bool is_oacc_serial + = (lookup_attribute ("oacc serial", + DECL_ATTRIBUTES (current_function_decl)) != NULL); + int fn_level = oacc_fn_attrib_level (attrs); + bool is_oacc_routine = (fn_level >= 0); + gcc_checking_assert (is_oacc_parallel + + is_oacc_kernels + + is_oacc_serial + + is_oacc_routine + == 1); + bool is_oacc_kernels_parallelized = (lookup_attribute ("oacc kernels parallelized", DECL_ATTRIBUTES (current_function_decl)) != NULL); + if (is_oacc_kernels_parallelized) + gcc_checking_assert (is_oacc_kernels); + + if (dump_file) + { + if (is_oacc_parallel) + fprintf (dump_file, "Function is OpenACC parallel offload\n"); + else if (is_oacc_kernels) + fprintf (dump_file, "Function is %s OpenACC kernels offload\n", + (is_oacc_kernels_parallelized + ? "parallelized" : "unparallelized")); + else if (is_oacc_serial) + fprintf (dump_file, "Function is OpenACC serial offload\n"); + else if (is_oacc_routine) + fprintf (dump_file, "Function is OpenACC routine level %d\n", + fn_level); + else + gcc_unreachable (); + } /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1 kernels, so remove the parallelism dimensions function attributes @@ -1559,22 +1592,10 @@ execute_oacc_device_lower () /* Discover, partition and process the loops. */ oacc_loop *loops = oacc_loop_discovery (); - int fn_level = oacc_fn_attrib_level (attrs); - - if (dump_file) - { - if (fn_level >= 0) - fprintf (dump_file, "Function is OpenACC routine level %d\n", - fn_level); - else if (is_oacc_kernels) - fprintf (dump_file, "Function is %s OpenACC kernels offload\n", - (is_oacc_kernels_parallelized - ? "parallelized" : "unparallelized")); - else - fprintf (dump_file, "Function is OpenACC parallel offload\n"); - } - unsigned outer_mask = fn_level >= 0 ? GOMP_DIM_MASK (fn_level) - 1 : 0; + unsigned outer_mask = 0; + if (is_oacc_routine) + outer_mask = GOMP_DIM_MASK (fn_level) - 1; unsigned used_mask = oacc_loop_partition (loops, outer_mask); /* OpenACC kernels constructs are special: they currently don't use the generic oacc_loop infrastructure and attribute/dimension processing. */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 66a6d133663..933d7664386 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -20,10 +20,10 @@ void PARALLEL () } /* Check the offloaded function's attributes. - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target entrypoint\\)\\)" 1 "ompexp" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc parallel, omp target entrypoint\\)\\)" 1 "ompexp" } } */ /* Check the offloaded function's classification and compute dimensions (will always be 1 x 1 x 1 for non-offloading compilation). { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccdevlow" } } { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } - { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c new file mode 100644 index 00000000000..94ace1b3c20 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c @@ -0,0 +1,29 @@ +/* Check offloaded function's attributes and classification for OpenACC + serial. */ + +/* { dg-additional-options "-O2" } + { dg-additional-options "-fopt-info-optimized-omp" } + { dg-additional-options "-fdump-tree-ompexp" } + { dg-additional-options "-fdump-tree-oaccdevlow" } */ + +#define N 1024 + +extern unsigned int *__restrict a; +extern unsigned int *__restrict b; +extern unsigned int *__restrict c; + +void SERIAL () +{ +#pragma acc serial loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */ + for (unsigned int i = 0; i < N; i++) + c[i] = a[i] + b[i]; +} + +/* Check the offloaded function's attributes. + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc serial, omp target entrypoint\\)\\)" 1 "ompexp" } } */ + +/* Check the offloaded function's classification and compute dimensions (will + always be 1 x 1 x 1 for non-offloading compilation). + { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccdevlow" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } + { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 index a23ea81609b..01f06bbcc27 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 @@ -22,10 +22,10 @@ program main end program main ! Check the offloaded function's attributes. -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target entrypoint\\)\\)" 1 "ompexp" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc parallel, omp target entrypoint\\)\\)" 1 "ompexp" } } ! Check the offloaded function's classification and compute dimensions (will ! always be 1 x 1 x 1 for non-offloading compilation). ! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccdevlow" } } ! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } -! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 new file mode 100644 index 00000000000..51061afd2c6 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 @@ -0,0 +1,31 @@ +! Check offloaded function's attributes and classification for OpenACC +! serial. + +! { dg-additional-options "-O2" } +! { dg-additional-options "-fopt-info-optimized-omp" } +! { dg-additional-options "-fdump-tree-ompexp" } +! { dg-additional-options "-fdump-tree-oaccdevlow" } + +program main + implicit none + integer, parameter :: n = 1024 + integer, dimension (0:n-1) :: a, b, c + integer :: i + + call setup(a, b) + + !$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } + do i = 0, n - 1 + c(i) = a(i) + b(i) + end do + !$acc end serial loop +end program main + +! Check the offloaded function's attributes. +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc serial, omp target entrypoint\\)\\)" 1 "ompexp" } } + +! Check the offloaded function's classification and compute dimensions (will +! always be 1 x 1 x 1 for non-offloading compilation). +! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } -- 2.17.1