From patchwork Mon Dec 22 17:41:11 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 423447 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 516921400D5 for ; Tue, 23 Dec 2014 04:41:34 +1100 (AEDT) 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=ELYXnzXZi6FH0KBQ IvKs1Kdd/CSMTFuP6YSpTgibggZGTAzgw4WdqFKMjCaT5FI8Xbar+2YnMvWOfRbK dsm7mWoSMoDi0BNUdf2E+xMmejTTl/S/4YdlbRQ+qRdnWGjn5pAkFgKVbgLLwJKE Qhwl8ojPJUWeNL6tofsYDrlCvx0= 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=uZyAEdG74GePm5Tw8ZA9Lc +Nrd0=; b=L6dmpycJdSW8a8/J2LFAvIGWr0RgqugfPJENYnFuqss1PyZQ7/ciap xwzZVj8OU53AKro+RHB94jiZsn4jbIWUWdbd0zOcobu5v8JqcsGzeOkAPPstaQLy brg4lFI8cHHyJDOEjTG5RQuq7aOa8rw1vzNbQYdSmzCwy5Z4qEdag= Received: (qmail 20137 invoked by alias); 22 Dec 2014 17:41:27 -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 20127 invoked by uid 89); 22 Dec 2014 17:41:26 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS 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; Mon, 22 Dec 2014 17:41:24 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Y36yu-0006IR-2n from Thomas_Schwinge@mentor.com ; Mon, 22 Dec 2014 09:41:20 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.181.6; Mon, 22 Dec 2014 17:41:18 +0000 From: Thomas Schwinge To: , Jakub Jelinek CC: Julian Brown Subject: Re: [PATCH 1/5] OpenACC 2.0 support for libgomp - OpenACC runtime, NVidia PTX/CUDA plugin (repost) In-Reply-To: <20141111135323.29e0f27b@octopus> References: <20140923191931.2177e60f@octopus> <20141111135323.29e0f27b@octopus> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (i586-pc-linux-gnu) Date: Mon, 22 Dec 2014 18:41:11 +0100 Message-ID: <871tnrzio8.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Tue, 11 Nov 2014 13:53:23 +0000, Julian Brown wrote: > On Tue, 23 Sep 2014 19:19:31 +0100 > Julian Brown wrote: > > > This patch contains the bulk of the OpenACC 2.0 runtime support, > > building around, or on top of, the OpenMP 4.0 support (as previously > > posted or already extant upstream) where we could. [...] > > Here is a new version of the OpenACC support patch for libgomp, rebased > on top of a version of Ilya Verbin's patches that I merged to a local > clone of trunk, and tested as far as possible without the > middle/front-end pieces, since those are not ready yet. > --- a/libgomp/target.c > +++ b/libgomp/target.c Code is split off of gomp_init_device into the new gomp_init_tables: > gomp_init_device (struct gomp_device_descr *devicep) > { > + /* Initialize the target device. */ > devicep->init_device_func (devicep->target_id); > + > + devicep->is_initialized = true; > +} > > +attribute_hidden void > +gomp_init_tables (const struct gomp_device_descr *devicep, > + struct gomp_memory_mapping *mm) > +{ > /* Get address mapping table for device. */ > struct mapping_table *table = NULL; > [...] ..., and a new function gomp_init_dev_tables added to call both of them: > +static void > +gomp_init_dev_tables (struct gomp_device_descr *devicep) > +{ > + gomp_init_device (devicep); > + gomp_init_tables (devicep, &devicep->mem_map); > +} ..., which is then used in GOMP_target: > @@ -673,7 +753,12 @@ GOMP_target (int device, void (*fn) (void *), const void *openmp_target, > + if (devicep != NULL && !devicep->is_initialized) > + gomp_init_dev_tables (devicep); ..., and GOMP_target_data: > @@ -724,7 +818,12 @@ GOMP_target_data (int device, const void *openmp_target, size_t mapnum, > + if (devicep != NULL && !devicep->is_initialized) > + gomp_init_dev_tables (devicep); ..., but not in GOMP_target_update: > @@ -771,15 +871,18 @@ GOMP_target_update (int device, const void *openmp_target, size_t mapnum, > + if (devicep != NULL && !devicep->is_initialized) > gomp_init_device (devicep); Committed to gomp-4_0-branch in r219023: commit 0020b7972a6f7f422bd9d4782021ff5b8e2b0f50 Author: tschwinge Date: Mon Dec 22 17:40:37 2014 +0000 libgomp: Restore code that initializes the address mapping tables in GOMP_target_update. libgomp/ * target.c (GOMP_target_update): To initialize, call gomp_init_dev_tables instead of gomp_init_device. With Intel MIC offloading (emulation), this fixes cases of "#pragma omp declare target" of variables, followed by "#pragma omp target update to([...])" before any "#pragma omp target" or "#pragma omp target data" is called: FAIL: libgomp.c/examples-4/e.53.3.c execution test FAIL: libgomp.c/examples-4/e.53.4.c execution test FAIL: libgomp.c/examples-4/e.53.5.c execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O0 execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O1 execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O2 execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O3 -fomit-frame-pointer execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O3 -fomit-frame-pointer -funroll-loops execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -O3 -g execution test FAIL: libgomp.fortran/examples-4/e.53.3.f90 -Os execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O0 execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O1 execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O2 execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O3 -fomit-frame-pointer execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O3 -fomit-frame-pointer -funroll-loops execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -O3 -g execution test FAIL: libgomp.fortran/examples-4/e.53.4.f90 -Os execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O0 execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O1 execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O2 execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O3 -fomit-frame-pointer execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O3 -fomit-frame-pointer -funroll-loops execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -O3 -g execution test FAIL: libgomp.fortran/examples-4/e.53.5.f90 -Os execution test git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@219023 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog.gomp | 3 +++ libgomp/target.c | 2 +- 2 files changed, 4 insertions(+), 1 deletion(-) Grüße, Thomas diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 26fdfe6..3aa9bf4 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-12-22 Thomas Schwinge + * target.c (GOMP_target_update): To initialize, call + gomp_init_dev_tables instead of gomp_init_device. + * target.c (gomp_map_vars) : Revert earlier changes. diff --git libgomp/target.c libgomp/target.c index 423bbee..8517a84 100644 --- libgomp/target.c +++ libgomp/target.c @@ -891,7 +891,7 @@ GOMP_target_update (int device, const void *offload_table, size_t mapnum, struct gomp_memory_mapping *mm = &devicep->mem_map; gomp_mutex_lock (&mm->lock); if (!devicep->is_initialized) - gomp_init_device (devicep); + gomp_init_dev_tables (devicep); gomp_mutex_unlock (&mm->lock); if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))