This is the mail archive of the
gcc@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4] Questions about "declare target" and "target update" pragmas
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: gcc at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>
- Date: Thu, 19 Mar 2015 17:49:47 +0300
- Subject: Re: [gomp4] Questions about "declare target" and "target update" pragmas
- Authentication-results: sourceware.org; auth=none
- References: <20140122155151 dot GA50489 at msticlxl57 dot ims dot intel dot com> <20150310165252 dot GC37666 at msticlxl57 dot ims dot intel dot com> <20150316184153 dot GA42550 at msticlxl57 dot ims dot intel dot com> <20150319134744 dot GW1746 at tucnak dot redhat dot com>
On Thu, Mar 19, 2015 at 14:47:44 +0100, Jakub Jelinek wrote:
> Here is untested patch. I'm going to check it in after bootstrap/regtest.
Thanks.
> > I am investigating run-fails on some benchmark, and have found a second
> > questionable place, where a function argument overrides a global array.
> > Just to be sure, is this a bug in the test?
> >
> > #pragma omp declare target
> > int a1[50], a2[50];
> > #pragma omp end declare target
> >
> > void foo (int a1[])
> > {
> > #pragma omp target
> > {
> > a1[10]++;
> > a2[10]++;
> > }
> > }
>
> That is a buggy test. int a1[] function argument is changed
> into int *a1, so it is actually
> #pragma omp target map(tofrom:a1, a2)
Actually, it copies only a1 pointer, since a2 points to the global array.
> {
> a1[10]++;
> a2[10]++;
> }
> which copies the a1 pointer to the device by value (no pointer
> transformation).
> Perhaps the testcase writer meant to use #pragma omp target map(a1[10])
> instead (or map(a1[0:50])?
If I understand correctly, it's not allowed to map global target arrays this
way, since it's already present in the initial device data environment:
2.9.4 declare target Directive
If a list item is a variable then the original variable is mapped to a corresponding
variable in the initial device data environment for all devices.
2.14.5 map Clause
If a corresponding list item of the original list item is in the enclosing device data
environment, the new device data environment uses the corresponding list item from the
enclosing device data environment. No additional storage is allocated in the new device
data environment and neither initialization nor assignment is performed, regardless of
the map-type that is specified.
So, to fix this testcase I can just remove the "int a1[]" function argument, and
add some "#pragma omp target update" where needed.
-- Ilya