From patchwork Wed Jun 1 21:35:44 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 628875 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 3rKkF20sFCz9t3n for ; Thu, 2 Jun 2016 07:36:05 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=Lbmlzhxt; 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 :subject:to:message-id:date:mime-version:content-type; q=dns; s= default; b=WWN71KJThmikM1K9q3ge3/IkCujM0tRdDoJhl22BEOVQufsXURy1Y a+ksSqQQSmY7NU8SV/Mg++LascN8X6oec5/VeHUp+i4JUAF6yKCrF8hp8Zqoo0qC Ox32DcHxcACJwnZPBnmIy6UATYCWXnsrOiCIDYbJRZAvL2VJih18nQ= 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 :subject:to:message-id:date:mime-version:content-type; s= default; bh=zfbR9/HAZeP0AZ0/Qm/DQ4gpTME=; b=LbmlzhxtSTGB4JWGf3Km EAruY+lkxQguXFijGv8WJQn43B9keIKpMDfKyewBvhNAM64lbxWGMBYbqz/eKbNv jkhdtYuaI+StWokA99orTJS6gthNML1oAVjeG/gZCtQfwxicyNbUagbmCvoN0xz3 Zy0Y+IVtoRexhrQK0O8gwz0= Received: (qmail 37110 invoked by alias); 1 Jun 2016 21:35:58 -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 37090 invoked by uid 89); 1 Jun 2016 21:35:57 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.2 spammy=teaches 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, 01 Jun 2016 21:35:47 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1b8Dnk-0003X2-Ry from Cesar_Philippidis@mentor.com ; Wed, 01 Jun 2016 14:35:44 -0700 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.224.2; Wed, 1 Jun 2016 14:35:44 -0700 From: Cesar Philippidis Subject: [PATCH] zero-length arrays in OpenACC To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek Message-ID: <574F5530.9080408@codesourcery.com> Date: Wed, 1 Jun 2016 14:35:44 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.8.0 MIME-Version: 1.0 This patch teaches c and c++ front ends and omp-low how to deal with subarray involving GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} data mappings. As the libgomp test case shows, it might be possible for a subarray to have zero length. This patch fixes that. I also updated *parser_oacc_declare not to handle GOMP_MAP_POINTER, because subarray pointers are now firstprivate. Is this OK for trunk? Cesar 2016-06-01 Cesar Philippidis gcc/c/ * c-parser.c (c_parser_oacc_declare): Don't scan for GOMP_MAP_POINTER. * c-typeck.c (handle_omp_array_sections): Mark data clauses with GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having zero-length subarrays. gcc/cp/ * parser.c (cp_parser_oacc_declare): Don't scan for GOMP_MAP_POINTER. * semantics.c (handle_omp_array_sections): Mark data clauses with GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having zero-length subarrays. gcc/ * omp-low.c (lower_omp_target): Mark data clauses with GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having zero-length subarrays. libgomp/ * testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c: New test. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index bca8653..9bd377c 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13622,11 +13622,6 @@ c_parser_oacc_declare (c_parser *parser) case GOMP_MAP_DEVICE_RESIDENT: break; - case GOMP_MAP_POINTER: - /* Generated by c_finish_omp_clauses from array sections; - avoid spurious diagnostics. */ - break; - case GOMP_MAP_LINK: if (!global_bindings_p () && (TREE_STATIC (decl) diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 1520c20..1e97749 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12439,6 +12439,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; break; default: diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 29a1b80..9568ef8 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35244,11 +35244,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) case GOMP_MAP_DEVICE_RESIDENT: break; - case GOMP_MAP_POINTER: - /* Generated by c_finish_omp_clauses from array sections; - avoid spurious diagnostics. */ - break; - case GOMP_MAP_LINK: if (!global_bindings_p () && (TREE_STATIC (decl) diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 8e682c5..d960747 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4995,6 +4995,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; break; default: diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 77bdb18..052e6d3 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -16208,6 +16208,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_RELEASE: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION; break; case GOMP_MAP_DELETE: diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c new file mode 100644 index 0000000..8954551 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c @@ -0,0 +1,45 @@ +/* Exercise zero-length sub-arrays. */ + +const int n = 10; + +void +subzero_present (int *a, int n) +{ +#pragma acc data present (a[0:n]) + ; +#pragma acc data pcopy (a[0:n]) + ; +#pragma acc data pcopyin (a[0:n]) + ; +#pragma acc data pcopyout (a[0:n]) + ; + +} + +void +subzero (int *a, int n) +{ +#pragma acc data create (a[0:n]) + ; +#pragma acc data copy (a[0:n]) + ; +#pragma acc data copyin (a[0:n]) + ; +#pragma acc data copyout (a[0:n]) + ; +} + +int +main () +{ + int a[n]; + +#pragma acc data copy (a[0:n]) + { + subzero_present (a, 0); + } + + subzero (a, 0); + + return 0; +}