From patchwork Wed Jul 27 08:59:02 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 653213 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3rzppJ1hTJz9t24 for ; Wed, 27 Jul 2016 18:59:33 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=ZGtv3tJD; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type :content-transfer-encoding; q=dns; s=default; b=qzQKEwPEP/ODhx0r ECnMeTUxyhG4C0ODZ9F8eaECGetAtK5XxhllVEa4HPGSGcTHI//eR9owAtTg9yIj 6sAwSQgcIj8yO8U9VaaQPup/CJGwoIu5s/ssVrkOxmrl1zyPknz2vfCh1d6kOR0E p2Gk7J7ma4utBw5czsPTrcV+vLI= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type :content-transfer-encoding; s=default; bh=qZQyQL5RSNvGl45b8b+NLB UKgcs=; b=ZGtv3tJDZdOjgiKMltdC8TAo/z40JWrYHBEe/yRq0yOOGkP3wK0IdM p46nymGcpn0/W4FfT7kIh9UQkNMREVJOmF0CJ5RPpETVZJrs4ByV5RfQb6gYUGyG /3y9VxKAjvCrKjIiG0XnB/sjN/GYBRxsOLkIs+f+mqf1WUnUsFFcE= Received: (qmail 119427 invoked by alias); 27 Jul 2016 08:59:24 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 118494 invoked by uid 89); 27 Jul 2016 08:59:23 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=entrypoint, SUCCESS X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Wed, 27 Jul 2016 08:59:13 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1bSKgI-0004U0-2C from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Wed, 27 Jul 2016 01:59:10 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Wed, 27 Jul 2016 09:59:08 +0100 From: Thomas Schwinge To: Subject: Test cases to check OpenACC offloaded function's attributes and classification User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Wed, 27 Jul 2016 10:59:02 +0200 Message-ID: <87zip3jw2x.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Hi! OK for trunk? commit 8200af082db5438be18bc60f721fcf21641c0d86 Author: Thomas Schwinge Date: Tue Jul 26 17:18:21 2016 +0200 Test cases to check OpenACC offloaded function's attributes and classification gcc/testsuite/ * c-c++-common/goacc/oaccdevlow-kernels.c: New file. * c-c++-common/goacc/oaccdevlow-parallel.c: Likewise. * c-c++-common/goacc/oaccdevlow-routine.c: Likewise. * gfortran.dg/goacc/oaccdevlow-kernels.f95: Likewise. * gfortran.dg/goacc/oaccdevlow-parallel.f95: Likewise. * gfortran.dg/goacc/oaccdevlow-routine.f95: Likewise. --- .../c-c++-common/goacc/oaccdevlow-kernels.c | 34 ++++++++++++++++++++ .../c-c++-common/goacc/oaccdevlow-parallel.c | 27 ++++++++++++++++ .../c-c++-common/goacc/oaccdevlow-routine.c | 29 +++++++++++++++++ .../gfortran.dg/goacc/oaccdevlow-kernels.f95 | 36 ++++++++++++++++++++++ .../gfortran.dg/goacc/oaccdevlow-parallel.f95 | 29 +++++++++++++++++ .../gfortran.dg/goacc/oaccdevlow-routine.f95 | 28 +++++++++++++++++ 6 files changed, 183 insertions(+) Grüße Thomas diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-kernels.c gcc/testsuite/c-c++-common/goacc/oaccdevlow-kernels.c new file mode 100644 index 0000000..14d650a --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-kernels.c @@ -0,0 +1,34 @@ +/* Check offloaded function's attributes and classification for OpenACC + kernels. */ + +/* { dg-additional-options "-O2" } + { dg-additional-options "-fdump-tree-ompexp" } + { dg-additional-options "-fdump-tree-parloops1-all" } + { dg-additional-options "-fdump-tree-oaccdevlow" } */ + +#define N (1024 * 512) + +extern unsigned int *__restrict a; +extern unsigned int *__restrict b; +extern unsigned int *__restrict c; + +void KERNELS () +{ +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) + 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__\\(\\(omp target entrypoint\\)\\)" 1 "ompexp" } } */ + +/* Check that exactly one OpenACC kernels loop is analyzed, and that it can be + parallelized. + { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops1" } } + { dg-final { scan-tree-dump-times "(?n)oacc function \\(0," 1 "parloops1" } } + { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */ + +/* Check the offloaded function's classification and compute dimensions (will + always be [1, 1, 1] for target compilation). + { dg-final { scan-tree-dump-times "(?n)Function is kernels offload" 1 "oaccdevlow" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } */ diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-parallel.c gcc/testsuite/c-c++-common/goacc/oaccdevlow-parallel.c new file mode 100644 index 0000000..63c372a --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-parallel.c @@ -0,0 +1,27 @@ +/* Check offloaded function's attributes and classification for OpenACC + parallel. */ + +/* { dg-additional-options "-O2" } + { dg-additional-options "-fdump-tree-ompexp" } + { dg-additional-options "-fdump-tree-oaccdevlow" } */ + +#define N (1024 * 512) + +extern unsigned int *__restrict a; +extern unsigned int *__restrict b; +extern unsigned int *__restrict c; + +void PARALLEL () +{ +#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) + 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__\\(\\(omp target entrypoint\\)\\)" 1 "ompexp" } } */ + +/* Check the offloaded function's classification and compute dimensions (will + always be [1, 1, 1] for target compilation). + { dg-final { scan-tree-dump-times "(?n)Function is parallel offload" 1 "oaccdevlow" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } */ diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c new file mode 100644 index 0000000..fa2eae7 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c @@ -0,0 +1,29 @@ +/* Check offloaded function's attributes and classification for OpenACC + routine. */ + +/* { dg-additional-options "-O2" } + { dg-additional-options "-fdump-tree-ompexp" } + { dg-additional-options "-fdump-tree-oaccdevlow" } */ + +#define N (1024 * 512) + +extern unsigned int *__restrict a; +extern unsigned int *__restrict b; +extern unsigned int *__restrict c; +#pragma acc declare copyin (a, b) create (c) + +#pragma acc routine worker +void ROUTINE () +{ +#pragma acc loop + 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__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */ + +/* Check the offloaded function's classification and compute dimensions (will + always be [1, 1, 1] for target compilation). + { dg-final { scan-tree-dump-times "(?n)Function is routine level 1" 1 "oaccdevlow" } } + { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } */ diff --git gcc/testsuite/gfortran.dg/goacc/oaccdevlow-kernels.f95 gcc/testsuite/gfortran.dg/goacc/oaccdevlow-kernels.f95 new file mode 100644 index 0000000..8ee641e --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/oaccdevlow-kernels.f95 @@ -0,0 +1,36 @@ +! Check offloaded function's attributes and classification for OpenACC +! kernels. + +! { dg-additional-options "-O2" } +! { dg-additional-options "-fdump-tree-ompexp" } +! { dg-additional-options "-fdump-tree-parloops1-all" } +! { 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 kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) + do i = 0, n - 1 + c(i) = a(i) + b(i) + end do + !$acc end kernels +end program main + +! Check the offloaded function's attributes. +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target entrypoint\\)\\)" 1 "ompexp" } } + +! Check that exactly one OpenACC kernels loop is analyzed, and that it can be +! parallelized. +! { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops1" } } +! { dg-final { scan-tree-dump-times "(?n)oacc function \\(0," 1 "parloops1" } } +! { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } + +! Check the offloaded function's classification and compute dimensions (will +! always be [1, 1, 1] for target compilation). +! { dg-final { scan-tree-dump-times "(?n)Function is kernels offload" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } diff --git gcc/testsuite/gfortran.dg/goacc/oaccdevlow-parallel.f95 gcc/testsuite/gfortran.dg/goacc/oaccdevlow-parallel.f95 new file mode 100644 index 0000000..0975eb8 --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/oaccdevlow-parallel.f95 @@ -0,0 +1,29 @@ +! Check offloaded function's attributes and classification for OpenACC +! parallel. + +! { dg-additional-options "-O2" } +! { 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 parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) + do i = 0, n - 1 + c(i) = a(i) + b(i) + end do + !$acc end parallel loop +end program main + +! Check the offloaded function's attributes. +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target entrypoint\\)\\)" 1 "ompexp" } } + +! Check the offloaded function's classification and compute dimensions (will +! always be [1, 1, 1] for target compilation). +! { dg-final { scan-tree-dump-times "(?n)Function is parallel offload" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } } diff --git gcc/testsuite/gfortran.dg/goacc/oaccdevlow-routine.f95 gcc/testsuite/gfortran.dg/goacc/oaccdevlow-routine.f95 new file mode 100644 index 0000000..a68b5eb --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/oaccdevlow-routine.f95 @@ -0,0 +1,28 @@ +! Check offloaded function's attributes and classification for OpenACC +! routine. + +! { dg-additional-options "-O2" } +! { dg-additional-options "-fdump-tree-ompexp" } +! { dg-additional-options "-fdump-tree-oaccdevlow" } + +subroutine ROUTINE + !$acc routine worker + integer, parameter :: n = 1024 + integer, dimension (0:n-1) :: a, b, c + integer :: i + + call setup(a, b) + + !$acc loop + do i = 0, n - 1 + c(i) = a(i) + b(i) + end do +end subroutine ROUTINE + +! Check the offloaded function's attributes. +! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 0, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } + +! Check the offloaded function's classification and compute dimensions (will +! always be [1, 1, 1] for target compilation). +! { dg-final { scan-tree-dump-times "(?n)Function is routine level 1" 1 "oaccdevlow" } } +! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }