Commit Message

Hi Bernd,
Here is updated patch for libgomp. It assumes that there is a constructor with
a call to GOMP_offload_register in every target image, created by mkoffload
tool. How does this look?
---
libgomp/libgomp.map | 1 +
libgomp/plugin-host.c | 58 ++++++++++++++++-
libgomp/target.c | 170 +++++++++++++++++++++++++++++++++++++++++++++----
3 files changed, 213 insertions(+), 16 deletions(-)

Comments

Hi,
On 03/08/2014 03:50 PM, Ilya Verbin wrote:
> Here is updated patch for libgomp. It assumes that there is a constructor with> a call to GOMP_offload_register in every target image, created by mkoffload> tool. How does this look?
LGTM. Shall I start committing my changes to the branch?
Bernd

2014-03-12 18:12 GMT+04:00 Bernd Schmidt <bernds@codesourcery.com>:
> LGTM. Shall I start committing my changes to the branch?
Yes, I think you should commit your changes.
And we will rewrite our part to use the new configure approach.
-- Ilya

On Mon, Mar 17, 2014 at 04:00:11PM +0100, Thomas Schwinge wrote:
> Hi!> > On Sat, 8 Mar 2014 18:50:15 +0400, Ilya Verbin <iverbin@gmail.com> wrote:> > --- a/libgomp/libgomp.map> > +++ b/libgomp/libgomp.map> > @@ -208,6 +208,7 @@ GOMP_3.0 {> > > > GOMP_4.0 {> > global:> > + GOMP_offload_register;> > GOMP_barrier_cancel;> > GOMP_cancel;> > GOMP_cancellation_point;> > Now that the GOMP_4.0 symbol version is being used in GCC trunk, and will> be in the GCC 4.9 release, can we still add new symbols to it here?> (Jakub?)
If GCC 4.9 release will not include that symbol, then it must be in a new
symbol version, e.g. GOMP_4.1 (note, the fact that GOMP_ symbol version
matched now the OpenMP standard version wasn't always true and might not be
true always either (or we could use GOMP_4.0.1 symver).
Jakub

On 17 Mar 16:00, Thomas Schwinge wrote:
> > GOMP_4.0 {> > global:> > + GOMP_offload_register;> > GOMP_barrier_cancel;> > GOMP_cancel;> > GOMP_cancellation_point;> > Now that the GOMP_4.0 symbol version is being used in GCC trunk, and will> be in the GCC 4.9 release, can we still add new symbols to it here?> (Jakub?)
I moved it to GOMP_4.0.1.
> > + /* This is the TYPE of device. */> > + int type;> > Use enum target_type instead of int?
Done.
> > + offload_images = realloc (offload_images,> > + (num_offload_images + 1)> > + * sizeof (struct offload_image_descr));> > +> > + if (offload_images == NULL)> > + return;> > Fail silently, or use gomp_realloc to fail loudly?
Replaced with gomp_realloc.
> > if (dir)> > closedir (dir);> > + free (offload_images);> > I suggest to set offload_images = NULL, for clarity.
Done.
> OK to commit, thanks!
Committed as r208657.
> Would it make sense to have device_run return a value to make it able to> indicate to libgomp that the function cannot be run on the device (for> whatever reason), and libgomp should use host-fallback execution?> (Probably that needs more thought and discussion, OK to defer.)
Consider the following example (using OpenMP, I don't know OpenACC :)
int foo ()
{
int x = 0;
/* offload_fn1 */
#pragma omp target map(to: x)
{
x += 5;
}
/* Some code on host without updating 'x' from target. */
/* offload_fn2 */
#pragma omp target map(from: x)
{
x += 10;
}
return x;
}
If both offload_fn1 and offload_fn2 are executed on host, everything is fine
and x = 15. The same goes to the case when both offload_fn1 and offload_fn2
are executed on target. But if offload_fn1 is executed on target and
offload_fn2 is executed on host, then 'x' will have incorrect value (10).
Therefore, I proposed to check for target device availability only during
initialization of the plugin. And to make a decision at this point, will
libgomp run all functions on host or on target. Probably libgomp should return
an error if something was executed on device, but then it becomes unavailable.
-- Ilya

On 03/12/2014 03:51 PM, Ilya Verbin wrote:
> 2014-03-12 18:12 GMT+04:00 Bernd Schmidt <bernds@codesourcery.com>:>> LGTM. Shall I start committing my changes to the branch?>> Yes, I think you should commit your changes.> And we will rewrite our part to use the new configure approach.
Done now. I think/hope that I've committed all the ones that are not
entirely related to ptx - let me know if you find there are any missing
pieces.
Bernd