From patchwork Wed Nov 11 17:19:55 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 542977 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 0E2871402B4 for ; Thu, 12 Nov 2015 04:20:08 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=f0/dp1Of; 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:to:cc :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=sJGnP3dzr1+RiBT3FeBAzdEttoNoD825mrnvNc1zwWNGpuL16/ MMOctTZ6nlZgZWgFYFihhWdJc9a7akwbZlO1KPfNAM8CJI1M/sMJ6qIr5E4Db5Ra 88YdfvZyKuuCMHAVLyXS12+He5AK36gpQaSWBf43yYvUVo6vphoi9BI4o= 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:to:cc :from:subject:message-id:date:mime-version:content-type; s= default; bh=G7VmTbtXMDdfwSf1gOsb21Fp194=; b=f0/dp1OfOJlu1nrWg19I O8jsxhFyh2ZEVmKZUfofG2MPPvQBJvbr74cvHR/vzdUZx89HOl/EzPpvRFmFT7cD xCzvJiYFPZDGoPDZCrfW1uC+u0ghmq0aedWsN8k4CcMPsWIHDB57R3JYc9rQDAAp pZxITE0Xa92NpIJdcZpUNuA= Received: (qmail 119083 invoked by alias); 11 Nov 2015 17:20:01 -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 118998 invoked by uid 89); 11 Nov 2015 17:20:00 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: Yes, score=6.1 required=5.0 tests=BAYES_40, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPAM_BODY, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-yk0-f172.google.com Received: from mail-yk0-f172.google.com (HELO mail-yk0-f172.google.com) (209.85.160.172) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Wed, 11 Nov 2015 17:19:59 +0000 Received: by ykdv3 with SMTP id v3so59856012ykd.0 for ; Wed, 11 Nov 2015 09:19:57 -0800 (PST) X-Received: by 10.13.199.130 with SMTP id j124mr10872731ywd.70.1447262397293; Wed, 11 Nov 2015 09:19:57 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id f133sm11626095ywa.27.2015.11.11.09.19.56 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Wed, 11 Nov 2015 09:19:56 -0800 (PST) To: Jakub Jelinek Cc: GCC Patches , Cesar Philippidis From: Nathan Sidwell Subject: open acc default data attribute Message-ID: <564378BB.8050400@acm.org> Date: Wed, 11 Nov 2015 12:19:55 -0500 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 Jakub, this patch implements default data attribute determination. The current behaviour defaults to 'copy' and ignores 'default(none)'. The patch corrects that. 1) We emit a diagnostic when 'default(none)' is in effect. The fortran FE emits some artificial decls that it doesn't otherwise annotate, which is why we check DECL_ARTIFICIAL. IIUC Cesar had a patch to address that but it needed some reworking? 2) 'copy' is the correct default for 'kernels' regions, but for a 'parallel' region, scalars should be 'firstprivate', which is what this patch implements. ok? nathan 2015-11-11 Nathan Sidwell gcc/ * gimplify.c (oacc_default_clause): New. (omp_notice_variable): Call it. gcc/testsuite/ * c-c++-common/goacc/data-default-1.c: New. libgomp/ * testsuite/libgomp.oacc-c-c++-common/default-1.c: New. Index: gcc/gimplify.c =================================================================== --- gcc/gimplify.c (revision 230169) +++ gcc/gimplify.c (working copy) @@ -5900,6 +5900,60 @@ omp_default_clause (struct gimplify_omp_ return flags; } + +/* Determine outer default flags for DECL mentioned in an OACC region + but not declared in an enclosing clause. */ + +static unsigned +oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) +{ + const char *rkind; + + switch (ctx->region_type) + { + default: + gcc_unreachable (); + + case ORT_ACC_KERNELS: + /* Everything under kernels are default 'present_or_copy'. */ + flags |= GOVD_MAP; + rkind = "kernels"; + break; + + case ORT_ACC_PARALLEL: + { + tree type = TREE_TYPE (decl); + + if (TREE_CODE (type) == REFERENCE_TYPE + || POINTER_TYPE_P (type)) + type = TREE_TYPE (type); + + if (AGGREGATE_TYPE_P (type)) + /* Aggregates default to 'present_or_copy'. */ + flags |= GOVD_MAP; + else + /* Scalars default to 'firstprivate'. */ + flags |= GOVD_FIRSTPRIVATE; + rkind = "parallel"; + } + break; + } + + if (DECL_ARTIFICIAL (decl)) + ; /* We can get compiler-generated decls, and should not complain + about them. */ + else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE) + { + error ("%qE not specified in enclosing OpenACC %s construct", + DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind); + error_at (ctx->location, "enclosing OpenACC %s construct", rkind); + } + else + gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED); + + return flags; +} + /* Record the fact that DECL was used within the OMP context CTX. IN_CODE is true when real code uses DECL, and false when we should merely emit default(none) errors. Return true if DECL is going to @@ -6023,7 +6077,12 @@ omp_notice_variable (struct gimplify_omp nflags |= GOVD_MAP | GOVD_EXPLICIT; } else if (nflags == flags) - nflags |= GOVD_MAP; + { + if ((ctx->region_type & ORT_ACC) != 0) + nflags = oacc_default_clause (ctx, decl, flags); + else + nflags |= GOVD_MAP; + } } found_outer: omp_add_variable (ctx, decl, nflags); Index: gcc/testsuite/c-c++-common/goacc/data-default-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/data-default-1.c (revision 0) +++ gcc/testsuite/c-c++-common/goacc/data-default-1.c (working copy) @@ -0,0 +1,37 @@ +/* { dg-do compile } */ + + +int main () +{ + int n = 2; + int ary[2]; + +#pragma acc parallel default (none) /* { dg-message "parallel construct" 2 } */ + { + ary[0] /* { dg-error "not specified in enclosing" } */ + = n; /* { dg-error "not specified in enclosing" } */ + } + +#pragma acc kernels default (none) /* { dg-message "kernels construct" 2 } */ + { + ary[0] /* { dg-error "not specified in enclosing" } */ + = n; /* { dg-error "not specified in enclosing" } */ + } + +#pragma acc data copy (ary, n) + { +#pragma acc parallel default (none) + { + ary[0] + = n; + } + +#pragma acc kernels default (none) + { + ary[0] + = n; + } + } + + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c (working copy) @@ -0,0 +1,87 @@ +/* { dg-do run } */ + +#include + +int test_parallel () +{ + int ok = 1; + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + + /* val defaults to firstprivate, ary defaults to copy. */ +#pragma acc parallel num_gangs (32) copy (ok) copy(ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop gang(static:1) + for (unsigned i = 0; i < 32; i++) + { + if (val != 2) + ok = 0; + val += i; + ary[i] = val; + } + } + + if (ondev) + { + if (!ok) + return 1; + if (val != 2) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} + +int test_kernels () +{ + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + + /* val defaults to copy, ary defaults to copy. */ +#pragma acc kernels copy(ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop + for (unsigned i = 0; i < 32; i++) + { + ary[i] = val; + val++; + } + } + + if (ondev) + { + if (val != 2 + 32) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} + +int main () +{ + if (test_parallel ()) + return 1; + + if (test_kernels ()) + return 1; + + return 0; +}