Re: Merge current set of OpenACC changes from gomp-4_0-branch

[Ilya Verbin <iverbin at gmail dot com>]

Hi!

On 15 Jan 21:20, Thomas Schwinge wrote:
> In r219682, I have committed to trunk our current set of OpenACC changes,
> which we had prepared on gomp-4_0-branch.  Thanks to everyone who has
> been contributing!

Unfortunately, it broke offloading from shared libraries (I mean common libs
with NEEDED entries, not dlopened).  Such things are not covered by the
testsuite, that's why you missed this issue.  Here is a simple testcase:

+++++ test.c: +++++

int f_aaa (void);

int main ()
{
  int x = f_aaa ();
  #pragma omp target
    x++;
  return x;
}

+++++ libaaa.c: +++++

int f_aaa (void)
{
  int x = 0;
  #pragma omp target
    x = 10;
  return x;
}

++++++++++

$ gcc -fopenmp -shared -fPIC libaaa.c -o libaaa.so
$ gcc -fopenmp -L. -laaa test.c
$ ./a.out
libgomp: Target function wasn't mapped


The problem seems to be here:

-gomp_register_images_for_device (struct gomp_device_descr *device)
+gomp_register_image_for_device (struct gomp_device_descr *device,
+				struct offload_image_descr *image)
 {
-  int i;
-  for (i = 0; i < num_offload_images; i++)
+  if (!device->offload_regions_registered
+      && (device->type == image->type
+	  || device->type == OFFLOAD_TARGET_TYPE_HOST))
     {
-      struct offload_image_descr *image = &offload_images[i];
-      if (image->type == device->type)
-	device->register_image_func (image->host_table, image->target_data);
+      device->register_image_func (image->host_table, image->target_data);
+      device->offload_regions_registered = true;
     }
 }

So, you don't assume that a device can have multiple images from multiple libs?

  -- Ilya