From patchwork Thu Dec 12 09:53:02 2013 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 300594 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)) (Client did not present a certificate) by ozlabs.org (Postfix) with ESMTPS id 63CD62C009F for ; Thu, 12 Dec 2013 20:53:21 +1100 (EST) 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:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=H3/eMiaWiqrJcxDy H4u+K6/3SWs61pSLRgptKxC1TTVUtIndFZte8zfoxNo5OtD4hPCgkT0WRcpiaq1p UIkIhTVNy4hbHONA7kmYstijNECGudZrhCJKeCbEaif6Inz5o41Ci69VKseny+63 C1MDLcSSEHu67AhYKVq0ae/25wA= 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:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=X+PIMOwF5JG/X1P7DIymLs eFcNI=; b=dSYn+IYH9wrfU1nlWgZwnWB62lt7EJcIzdHubahZ6Gv84LEAhWEnM1 ghIiZ5ncSCxsltBOranNxfat6TIXGPWBlIuInB41H5yHFnC5hY2zr9Q0Sh6WzXLL 5GPFBlXjSZG85ztPnQy8UeqPh0eVG0PZGR4byEEp/brAr9lJfQVKo= Received: (qmail 32739 invoked by alias); 12 Dec 2013 09:53:15 -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 32730 invoked by uid 89); 12 Dec 2013 09:53:14 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 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 ESMTP; Thu, 12 Dec 2013 09:53:13 +0000 Received: from svr-orw-fem-01.mgc.mentorg.com ([147.34.98.93]) by relay1.mentorg.com with esmtp id 1Vr2xA-0007lo-T0 from Thomas_Schwinge@mentor.com ; Thu, 12 Dec 2013 01:53:08 -0800 Received: from SVR-IES-FEM-02.mgc.mentorg.com ([137.202.0.106]) by svr-orw-fem-01.mgc.mentorg.com over TLS secured channel with Microsoft SMTPSVC(6.0.3790.4675); Thu, 12 Dec 2013 01:53:08 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.2.247.3; Thu, 12 Dec 2013 09:53:06 +0000 From: Thomas Schwinge To: Jakub Jelinek CC: , Richard Henderson , "Michael V. Zolotukhin" , Subject: GOMP_target: alignment (was: [gomp4] #pragma omp target* fixes) In-Reply-To: <20130905161105.GL23437@tucnak.redhat.com> References: <20130905161105.GL23437@tucnak.redhat.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (i486-pc-linux-gnu) Date: Thu, 12 Dec 2013 10:53:02 +0100 Message-ID: <87zjo6s8e9.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Thu, 5 Sep 2013 18:11:05 +0200, Jakub Jelinek wrote: > 3) I figured out we need to tell the runtime library not just > address, size and kind, but also alignment (we won't need that for > the #pragma omp declare target global vars though), so that the > runtime library can properly align it. As TYPE_ALIGN/DECL_ALIGN > is in bits and is 32 bit wide, when that is in bytes and we only care > about power of twos, I've decided to encode it in the upper 5 bits > of the kind (lower 3 bits are used for OMP_CLAUSE_MAP_* kind). Unfortunately, this scheme breaks down with OpenACC: we need an additional bit to codify a flag for present_or_* map clauses (meaning: only map the data (allocate/to/from/tofrom, as for OpenMP) if not already present on the device). With five bits available for the OpenMP case, we can describe alignments up to 2 GiB, and I've empirically found on my development system that the largest possible alignment is MAX_OFILE_ALIGNMENT, 256 MiB for ELF systems, so that's fine. But with only four bits available, we get to describe alignments up to 1 << ((1 << 4) - 1) = 32 KiB, which is too small -- even though it'd be fine for "normal" usage of __attribute__ ((aligned (x))). So it seems our options are to use a bigger datatype for the kinds array, to split off from the kinds array a new alignments array, or to generally switch to using an array of a struct containing hostaddr, size, alignment, kind. The latter would require additional changes in the child_fn. As it's an ABI change no matter what, would you like to see this limited to OpenACC? Changing it also for OpenMP's GOMP_target would have the advantage to have them not diverge (especially at the generating side in omp-low.c's lowering functions), but I'm not sure whether such an ABI change would easily be possible now, with the OpenMP 4 support merged into trunk -- though, it is not yet part of a regular GCC release? > --- gcc/omp-low.c.jj 2013-09-05 09:19:03.000000000 +0200 > +++ gcc/omp-low.c 2013-09-05 17:11:14.693638660 +0200 > @@ -9342,6 +9349,11 @@ lower_omp_target (gimple_stmt_iterator * | unsigned char tkind = 0; | switch (OMP_CLAUSE_CODE (c)) | { | case OMP_CLAUSE_MAP: | tkind = OMP_CLAUSE_MAP_KIND (c); | break; | case OMP_CLAUSE_TO: | tkind = OMP_CLAUSE_MAP_TO; | break; | case OMP_CLAUSE_FROM: | tkind = OMP_CLAUSE_MAP_FROM; | break; > default: > gcc_unreachable (); > } > + unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); > + if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) > + talign = DECL_ALIGN_UNIT (ovar); > + talign = ceil_log2 (talign); > + tkind |= talign << 3; > CONSTRUCTOR_APPEND_ELT (vkind, purpose, > build_int_cst (unsigned_char_type_node, > tkind)); The use of OMP_CLAUSE_MAP_* on the generating and integer numerals on the receiving (libgomp) side is a bit unesthetic, likewise for the hard-coded 3 in the bit shift. What would be the standard GCC way of sharing a description of the tkind layout between gcc/omp-low.c and libgomp/target.c? Are we allowed to #include (a new header file) libgomp/target.h from gcc/omp-low.c? To avoid silent breakage should alignments bigger than 2 GiB be allowed in a distant future, would a check like the following be appropriate? Grüße, Thomas --- gcc/omp-low.c +++ gcc/omp-low.c @@ -10378,6 +10383,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); + const unsigned int talign_max + = 1 << ((1 << (BITS_PER_UNIT - 3)) - 1); + if (talign > talign_max) + sorry ("can't encode alignment of %u bytes, which is bigger than " + "%u bytes", talign, talign_max); talign = ceil_log2 (talign); tkind |= talign << 3; CONSTRUCTOR_APPEND_ELT (vkind, purpose,