From patchwork Thu Nov 26 09:50:46 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 548984 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 5C50F1402DD for ; Thu, 26 Nov 2015 20:51:03 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=BKHdlxHY; 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:date :from:to:cc:subject:message-id:reply-to:references:mime-version :content-type:in-reply-to; q=dns; s=default; b=HRT0G6jUG7VIsd37p rWsMHOQEe9U9CN/8qfDzP0dHJYY73rpXc28LRpQLX9H1qUHbaGRPz8zDX8/qL8uK EyqtwQkVE8oZGTnVQEFDUHpjFv6dvdALEUhWUKF618Em4FTTP0miSQQpQbvVBYYG Z45me+XIXmbOs1xXHR46i8o0Lc= 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:date :from:to:cc:subject:message-id:reply-to:references:mime-version :content-type:in-reply-to; s=default; bh=s/zY4RbxRYOH15VYU2iqJEA 67rQ=; b=BKHdlxHYDwpJR18Bc+CaN0uziETl4leybNgCBzlmNaTqRJugc9StKwE to03DhoY7zC7TWhMScNZ1HSKm3+YYInUWN5p79BdY+znDKE9As874sAvvcmZjg2M PLjl7oUweBRXcenDR7vsXRcGuBKSrM8T2mHZs/rRm65SNFzuaZVE= 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-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 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 10:50:46 +0100 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-Disposition: inline In-Reply-To: <20151110103936.GX5675@tucnak.redhat.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes 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. Jakub --- 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; +}