This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: Partial Offloading (was: [hsa merge 07/10] IPA-HSA pass)
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Thomas Schwinge <thomas at codesourcery dot com>, Jakub Jelinek <jakub at redhat dot com>
- Cc: Martin Jambor <mjambor at suse dot cz>, Alexander Monakov <amonakov at ispras dot ru>, gcc-patches at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>
- Date: Wed, 17 Feb 2016 19:41:15 +0300
- Subject: Re: Partial Offloading (was: [hsa merge 07/10] IPA-HSA pass)
- Authentication-results: sourceware.org; auth=none
- References: <20160115145323 dot GL3982 at virgil dot suse dot cz> <20160115150149 dot GX3017 at tucnak dot redhat dot com> <20160115160234 dot GO3982 at virgil dot suse dot cz> <20160115160954 dot GZ3017 at tucnak dot redhat dot com> <20160115163814 dot GB48907 at msticlxl57 dot ims dot intel dot com> <20160115164522 dot GA3017 at tucnak dot redhat dot com> <20160115180547 dot GC48907 at msticlxl57 dot ims dot intel dot com> <20160120185330 dot GC49431 at msticlxl57 dot ims dot intel dot com> <20160122101941 dot GP3017 at tucnak dot redhat dot com> <871t92x7f0 dot fsf at kepler dot schwinge dot homeip dot net>
On Thu, Jan 28, 2016 at 12:36:19 +0100, Thomas Schwinge wrote:
> I made an attempt to capture the recent discussion (plus my own
> ideas/understanding) in this new section:
> <https://gcc.gnu.org/wiki/Offloading#Partial_Offloading>. Please
> change/extend, as required.
Thanks for summarizing this.
I'm not very happy how -foffload=disable works in GCC 6, here is a testcase:
int main ()
{
int x = 10;
#pragma omp target data map (from: x)
#pragma omp target map (alloc: x)
x = 20;
if (x != 10 && x != 20)
__builtin_abort ();
}
On the system with non-shared accelerator it will abort, because "#pragma omp
target data" behaves like offloading is enabled, but "#pragma omp target" runs
on the host. As the result, at the end of the *target data* region, it tries to
receive x from target and receives 0, or crashes.
We can forbid -foffload=disable option, but I think it's very useful, e.g. for
comparing performance of host vs. accelerator using the same compiler, etc.
Or if the system contains 2 different accelerators, someone might want to
compile only for the first, but libgomp will load 2 plugins, and the program
will crash (instead of doing fallback) if it will try to use the second device.
So, maybe we still need something like this patch?
https://gcc.gnu.org/ml/gcc-patches/2015-04/msg01033.html
-- Ilya