From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 25865 invoked by alias); 26 Nov 2015 09:50:56 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 25847 invoked by uid 89); 26 Nov 2015 09:50:55 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.3 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Thu, 26 Nov 2015 09:50:54 +0000 Received: from int-mx11.intmail.prod.int.phx2.redhat.com (int-mx11.intmail.prod.int.phx2.redhat.com [10.5.11.24]) by mx1.redhat.com (Postfix) with ESMTPS id 5DDA9A1702; Thu, 26 Nov 2015 09:50:53 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-34.ams2.redhat.com [10.36.116.34]) by int-mx11.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id tAQ9ooZY010238 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Thu, 26 Nov 2015 04:50:52 -0500 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id tAQ9ommx001382; Thu, 26 Nov 2015 10:50:49 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id tAQ9okFC001381; Thu, 26 Nov 2015 10:50:46 +0100 Date: Thu, 26 Nov 2015 09:51:00 -0000 From: Jakub Jelinek To: Alexander Monakov Cc: gcc-patches@gcc.gnu.org, Dmitry Melnik Subject: Re: [gomp4 06/14] omp-low: copy omp_data_o to shared memory on NVPTX Message-ID: <20151126095046.GD5675@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <1445366076-16082-1-git-send-email-amonakov@ispras.ru> <1445366076-16082-7-git-send-email-amonakov@ispras.ru> <20151110103936.GX5675@tucnak.redhat.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20151110103936.GX5675@tucnak.redhat.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-11/txt/msg03191.txt.bz2 On Tue, Nov 10, 2015 at 11:39:36AM +0100, Jakub Jelinek wrote: > On Tue, Nov 03, 2015 at 05:25:53PM +0300, Alexander Monakov wrote: > > Here's an alternative patch that does not depend on exposure of shared-memory > > address space, and does not try to use pass_late_lower_omp. It's based on > > Bernd's suggestion to transform > > FYI, I've committed a new testcase to gomp-4_5-branch that covers various > target data sharing/team sharing/privatization parallel > sharing/privatization offloading cases. And another testcase, this time using only OpenMP 4.0 features, and trying to test the behavior of addressable vars in declare target functions where it is not clear if they are executed in teams, distribute or parallel for contexts. Wanted to look what LLVM generates here (tried llvm trunk), but they are unable to parse #pragma omp distribute or #pragma omp declare target, so it is hard to guess anything. Tested with XeonPhi offloading as well as host fallback, committed to trunk. 2015-11-26 Jakub Jelinek * testsuite/libgomp.c/target-35.c: New test. --- libgomp/testsuite/libgomp.c/target-35.c (revision 0) +++ libgomp/testsuite/libgomp.c/target-35.c (working copy) @@ -0,0 +1,129 @@ +#include +#include + +#pragma omp declare target +__attribute__((noinline)) +void +foo (int x, int y, int z, int *a, int *b) +{ + if (x == 0) + { + int i, j; + for (i = 0; i < 64; i++) + #pragma omp parallel for shared (a, b) + for (j = 0; j < 32; j++) + foo (3, i, j, a, b); + } + else if (x == 1) + { + int i, j; + #pragma omp distribute dist_schedule (static, 1) + for (i = 0; i < 64; i++) + #pragma omp parallel for shared (a, b) + for (j = 0; j < 32; j++) + foo (3, i, j, a, b); + } + else if (x == 2) + { + int j; + #pragma omp parallel for shared (a, b) + for (j = 0; j < 32; j++) + foo (3, y, j, a, b); + } + else + { + #pragma omp atomic + b[y] += z; + #pragma omp atomic + *a += 1; + } +} + +__attribute__((noinline)) +int +bar (int x, int y, int z) +{ + int a, b[64], i; + a = 8; + for (i = 0; i < 64; i++) + b[i] = i; + foo (x, y, z, &a, b); + if (x == 0) + { + if (a != 8 + 64 * 32) + return 1; + for (i = 0; i < 64; i++) + if (b[i] != i + 31 * 32 / 2) + return 1; + } + else if (x == 1) + { + int c = omp_get_num_teams (); + int d = omp_get_team_num (); + int e = d; + int f = 0; + for (i = 0; i < 64; i++) + if (i == e) + { + if (b[i] != i + 31 * 32 / 2) + return 1; + f++; + e = e + c; + } + else if (b[i] != i) + return 1; + if (a < 8 || a > 8 + f * 32) + return 1; + } + else if (x == 2) + { + if (a != 8 + 32) + return 1; + for (i = 0; i < 64; i++) + if (b[i] != i + (i == y ? 31 * 32 / 2 : 0)) + return 1; + } + else if (x == 3) + { + if (a != 8 + 1) + return 1; + for (i = 0; i < 64; i++) + if (b[i] != i + (i == y ? z : 0)) + return 1; + } + return 0; +} +#pragma omp end declare target + +int +main () +{ + int i, j, err = 0; + #pragma omp target map(tofrom:err) + #pragma omp teams reduction(+:err) + err += bar (0, 0, 0); + if (err) + abort (); + #pragma omp target map(tofrom:err) + #pragma omp teams reduction(+:err) + err += bar (1, 0, 0); + if (err) + abort (); + #pragma omp target map(tofrom:err) + #pragma omp teams reduction(+:err) + #pragma omp distribute + for (i = 0; i < 64; i++) + err += bar (2, i, 0); + if (err) + abort (); + #pragma omp target map(tofrom:err) + #pragma omp teams reduction(+:err) + #pragma omp distribute + for (i = 0; i < 64; i++) + #pragma omp parallel for reduction(+:err) + for (j = 0; j < 32; j++) + err += bar (3, i, j); + if (err) + abort (); + return 0; +} Jakub