This is the mail archive of the
gcc@gcc.gnu.org
mailing list for the GCC project.
Re: [RFC] Offloading Support in libgomp
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: "Michael V. Zolotukhin" <michael dot v dot zolotukhin at gmail dot com>, Kirill Yukhin <kirill dot yukhin at gmail dot com>, Richard Henderson <rth at redhat dot com>, GCC Development <gcc at gcc dot gnu dot org>, triegel at redhat dot com, Sergey Ostanevich <sergos dot gnu at gmail dot com>
- Date: Fri, 31 Jan 2014 21:18:33 +0400
- Subject: Re: [RFC] Offloading Support in libgomp
- Authentication-results: sourceware.org; auth=none
- References: <20130826132936 dot GB40923 at msticlxl57 dot ims dot intel dot com> <20130826141117 dot GF21876 at tucnak dot zalov dot cz> <20130827112609 dot GA4093 at msticlxl57 dot ims dot intel dot com> <20130827113956 dot GH21876 at tucnak dot zalov dot cz> <20130827115538 dot GB4093 at msticlxl57 dot ims dot intel dot com> <20130913093417 dot GA30181 at msticlxl57 dot ims dot intel dot com> <20130916093549 dot GP1817 at tucnak dot redhat dot com> <20130917120454 dot GE60139 at msticlxl57 dot ims dot intel dot com> <20130917122927 dot GY1817 at tucnak dot redhat dot com> <20131028104110 dot GA11997 at msticlxl57 dot ims dot intel dot com> <20131029070331 dot GT30970 at tucnak dot zalov dot cz>
Looks like there is a bug (in GOMP_target lowering? or in
gomp_map_vars_existing?)
The reproducer:
#define N 1000
void foo ()
{
int *a = malloc (N * sizeof (int));
printf ("1: %p\n", a);
#pragma omp target data map(tofrom: a[0:N])
{
printf ("2: %p\n", a);
#pragma omp target
{
int i;
for (i = 0; i < N; i++)
a[i] = i;
}
printf ("3: %p\n", a);
}
printf ("4: %p\n", a);
free (a);
}
Here GOMP_target believes that the pointer 'a' has a type TOFROM, so
it sets copy_from to true for the existing mapping of the pointer 'a',
that was mapped in GOMP_target_data. Therefore the output is
incorrect:
1: [host addr]
2: [host addr]
3: [host addr]
4: [target addr]
-- Ilya