From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 4017 invoked by alias); 28 Aug 2013 11:38:05 -0000 Mailing-List: contact gcc-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-owner@gcc.gnu.org Received: (qmail 4008 invoked by uid 89); 28 Aug 2013 11:38:05 -0000 Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 28 Aug 2013 11:38:05 +0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-4.5 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from int-mx01.intmail.prod.int.phx2.redhat.com (int-mx01.intmail.prod.int.phx2.redhat.com [10.5.11.11]) by mx1.redhat.com (8.14.4/8.14.4) with ESMTP id r7SBc2D3022856 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=OK); Wed, 28 Aug 2013 07:38:02 -0400 Received: from tucnak.zalov.cz (vpn1-5-98.ams2.redhat.com [10.36.5.98]) by int-mx01.intmail.prod.int.phx2.redhat.com (8.13.8/8.13.8) with ESMTP id r7SBc0sR006465 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-SHA bits=256 verify=NO); Wed, 28 Aug 2013 07:38:02 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.7/8.14.7) with ESMTP id r7SBc0Sj024945; Wed, 28 Aug 2013 13:38:00 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.7/8.14.7/Submit) id r7SBbxqT024944; Wed, 28 Aug 2013 13:37:59 +0200 Date: Wed, 28 Aug 2013 17:15:00 -0000 From: Jakub Jelinek To: Richard Biener Cc: "Michael V. Zolotukhin" , Kirill Yukhin , Richard Henderson , GCC Development , Torvald Riegel Subject: Re: [RFC] Offloading Support in libgomp Message-ID: <20130828113759.GU21876@tucnak.zalov.cz> Reply-To: Jakub Jelinek References: <20130822140810.GA27868@msticlxl57.ims.intel.com> <20130822142814.GB1814@tucnak.redhat.com> <20130823092810.GA36483@msticlxl57.ims.intel.com> <20130823095250.GJ1814@tucnak.redhat.com> <20130823153052.GA2974@msticlxl57.ims.intel.com> <20130823161631.GO1814@tucnak.redhat.com> <20130826115911.GA40923@msticlxl57.ims.intel.com> <20130828110601.GS21876@tucnak.zalov.cz> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: User-Agent: Mutt/1.5.21 (2010-09-15) X-SW-Source: 2013-08/txt/msg00333.txt.bz2 On Wed, Aug 28, 2013 at 01:21:53PM +0200, Richard Biener wrote: > My thought was that we need to have control over scheduling and thus have > a single runtime to be able to execute the following in parallel on the > accelerator and the CPU: > > #pragma omp parallel > { > #pragma omp target > for (;;) > ... > #pragma omp for > for (;;) > ... > } > #pragma omp wait > > that is, the omp target dispatch may not block the CPU. I can hardly OpenMP #pragma omp target blocks the host CPU until the accelerator code finishes. So if the goal is to spawn some accelerator code in parallel with parallelized host code, you'd need to make the code more complicated. I guess you could #pragma omp parallel { #pragma omp single #pragma omp target { #pragma omp parallel ... } #pragma omp for schedule(dynamic, N) for (;;) ... } or similar, then only one of the host parallel threads would spawn the target code, wait for it to be done and other threads in the mean time would do the worksharing (and the dynamic schedule would make sure that if the target region took long time, then no work or almost no work would be scheduled for the thread executing the target region). > > In the Intel MIC case (the only thing I've looked briefly at for how the > > offloading works - the COI library) you can load binaries and shared > > libraries either from files or from host memory image, so e.g. you can > > embed the libgomp library, some kind of libm and some kind of libc > > (would that be glibc, newlib, something else?) compiled for the target > > into some data section inside of the plugin or something > > (or load it from files of course). No idea how you do this in the > > HSAIL case, or PTX. > > For HSA you can do arbitrary calls to CPU code (that will then of course > execute on the CPU). GCC compiles into assembly or bytecode for HSAIL, right, and that then is further processed by some (right now proprietary?) blob. The question is does this allow linking of multiple HSAIL bytecode objects/libraries, etc. Say you have something providing (a subset of) C library, math library, libgomp, then say for OpenMP one host shared library provides some #pragma omp declare target ... #pragma omp end declare target routine, and another shared library uses #pragma omp target and calls that routine from there. So, I'd assume you have some HSAIL assembly/bytecode in each of the shared libraries, can you link that together and tell the runtime to execute some (named?) routine in there? Jakub