On Fri, Jan 15, 2016 at 17:45:22 +0100, Jakub Jelinek wrote: > On Fri, Jan 15, 2016 at 07:38:14PM +0300, Ilya Verbin wrote: > > On Fri, Jan 15, 2016 at 17:09:54 +0100, Jakub Jelinek wrote: > > > On Fri, Jan 15, 2016 at 05:02:34PM +0100, Martin Jambor wrote: > > > > How do other accelerators cope with the situation when half of the > > > > application is compiled with the accelerator disabled? (Would some of > > > > their calls to GOMP_target_ext lead to abort?) > > > > > > GOMP_target_ext should never abort (unless internal error), worst case it > > > just falls back into the host fallback. > > > > Wouldn't that lead to hard-to-find problems in case of nonshared memory? > > I mean when someone expects that all target regions are executed on the > > device, > > but in fact some of them are silently executed on the host with different > > data > > environment. > > E.g. for HSA it really shouldn't matter, as it is shared memory accelerator. > For XeonPhi we hopefully can offload anything.
As you said, if compilation of target image fails with ICE or somehow, host fallback and offloading to other targets should still work: https://gcc.gnu.org/ml/gcc-patches/2015-02/msg00951.html That patch was not applied, but it can be simulated by -foffload=disable, I've created a testcase: $ cat main.c #pragma omp declare target int x; #pragma omp end declare target extern int foo (); int main () { int shared_mem = 0; #pragma omp target map (alloc: x, shared_mem) { x = 10; shared_mem = 1; } x = 20; int r = foo (); if (!shared_mem && r != 100) __builtin_abort (); return 0; } $ cat liba.c #pragma omp declare target extern int x; #pragma omp end declare target int foo () { int r; #pragma omp target map (from: r) map (alloc: x) r = x * x; return r; } $ gcc -fopenmp -fPIC -shared liba.c -o liba.so -foffload=disable $ gcc -fopenmp -L. -la main.c Currently it prints "libgomp: Target function wasn't mapped", but after this change: --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1390,7 +1390,7 @@ gomp_get_target_fn_addr (struct gomp_device_descr *devicep, splay_tree_key tgt_fn = splay_tree_lookup (&devicep->mem_map, &k); gomp_mutex_unlock (&devicep->lock); if (tgt_fn == NULL) - gomp_fatal ("Target function wasn't mapped"); + return NULL; ... it will fail at __builtin_abort, but without -foffload=disable it will pass. -- Ilya