Re: [RFC][gomp4] Offloading patches (1/3): Add '-fopenmp_target' option

[Andrey Turetskiy <andrey dot turetskiy at gmail dot com>]

Hi Bernd,

We have some testcases, but they require XeonPhi hardware and a
working libgomp plugin. Our current version of the plugin depends on
some libraries, that are not open-sourced yet, so currently we can’t
share it.

However, you could examine what these patches do, making the following steps:
1) Build GCC with patches:
        http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01484.html
        http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01485.html
        http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01486.html
        http://gcc.gnu.org/ml/gcc-patches/2013-12/msg01896.html
2) Set environment variables (e.g. for two ‘targets’):
        export OFFLOAD_TARGET_NAMES=mic:hsail              (for now
names don’t really matter)
        export OFFLOAD_TARGET_COMPILERS=./gcc:./gcc    (use GCC with
patches above as target compiler, because it must support the
-fopenmp_target option)
3) Build any example with #pragma omp target (e.g. see attachment):
        ./gcc -flto -fopenmp test.c -o test.exe
    Options -flto and -fopenmp are necessary for using.

Now you have a binary with target images embedded and tables properly
filled. You can’t run it due to reasons mentioned above, though you
could examine it with objdump/nm/readelf to see new sections and their
content: there will be .offload_image_section with ‘target’ code and
.offload_func_table_section with ‘target’ function table.

On Tue, Jan 21, 2014 at 6:40 PM, Bernd Schmidt <bernds@codesourcery.com> wrote:
> On 12/17/2013 12:35 PM, Michael V. Zolotukhin wrote:
>>
>> Here is a set of patches implementing one more piece of offloading support
>> in
>> GCC.  These three patches allow to build a host binary with target image
>> and all
>> tables embedded.  Along with patches for libgomp and libgomp plugin, which
>> hopefully will be sent soon, that gives a functional and runnable
>> executable (or
>> DSO) with actual offloading to MIC.
>
>
> Do you have a testcase that can be used to see what this does in action?
>
>
> Bernd
>



-- 
Best regards,
Andrey Turetskiy

#define NUM 128

#pragma omp declare target
int cnt = 100;

int calc (int x)
{
  return cnt++ * x;
}
#pragma omp end declare target

int main ()
{
  int a[NUM];
  int i;

  cnt = 0;

  #pragma omp target update to(cnt)

  a[0] = 1;
  for (i = 1; i < NUM; i++)
    a[i] = -a[i-1];

  #pragma omp target data map(to: a)
    {
      for (i = 0; i < NUM; i++)
        a[i] = 2*a[i];

      #pragma omp target
        {
          for (i = 0; i < NUM; i++)
            a[i] = calc (a[i]);
        }

      #pragma omp target update from(a)
    }

  return 0;
}