On 11/09/15 08:46, Jakub Jelinek wrote:
On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:
Index: gcc/gimplify.c
===================================================================
If you want to switch to hexadecimal, you should change all values
in the enum to hexadecimal for consistency.
ok.
/* Gimplify hashtable helper. */
@@ -377,6 +383,12 @@ new_omp_context (enum omp_region_type re
else
c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+ c->combined_loop = false;
+ c->distribute = false;
+ c->target_map_scalars_firstprivate = false;
+ c->target_map_pointers_as_0len_arrays = false;
+ c->target_firstprivatize_array_bases = false;
Why this? c is XCNEW allocated, so zero initialized.
I presumed it necessary, as it was on the branch. will remove.
@@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct
/* We shouldn't be re-adding the decl with the same data
sharing class. */
gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
- /* The only combination of data sharing classes we should see is
- FIRSTPRIVATE and LASTPRIVATE. */
nflags = n->value | flags;
- gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
- == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
+ /* The only combination of data sharing classes we should see is
+ FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
+ reduction variables to be used in data sharing clauses. */
+ gcc_assert ((ctx->region_type & ORT_ACC) != 0
+ || ((nflags & GOVD_DATA_SHARE_CLASS)
+ == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
|| (flags & GOVD_DATA_SHARE_CLASS) == 0);
Are you sure you want to give up on any kind of consistency checks for
OpenACC? If only reduction is special on OpenACC, perhaps you could tweak
the assert for that instead? Something that can be done incrementally of
course.
Will investigate (later)
+
+ /* OpenMP doesn't look in outer contexts to find an
+ enclosing data clause. */
I'm puzzled by the comment. OpenMP does look in outer context for clauses
that need that (pretty much all closes but private), that is the do_outer:
recursion in omp_notice_variable. Say for firstprivate in order to copy (or
copy construct) the private variable one needs the access to the outer
context's var etc.).
So perhaps it would help to document what you are doing here for OpenACC and
why.
Ok. It seemed (and it may become clearer with default handling added), that
OpenACC and OpenMP scanned scopes in opposite orders. I remember trying to
get the ACC code to scan in the same order, but came up blank. Anyway, you're
right, it should say what OpenACC is trying.
The main issue I have is with the omp-low.c changes.
I see:
"2.5.9
private clause
The private clause is allowed on the parallel construct; it declares that a
copy of each
item on the list will be created for each parallel gang.
2.5.10
firstprivate clause
The firstprivate clause is allowed on the parallel construct; it declares that
a copy
of each item on the list will be created for each parallel gang, and that the
copy will be
initialized with the value of that item on the host when the parallel construct
is
encountered."
but looking at what you actually emit looks like standard present_copyin
clause I think with a private variable defined in the region where the
value of the present_copyin mapped variable is assigned to the private one.
This I'm afraid performs often two copies rather than just one (one to copy
the host value to the present_copyin mapped value, another one in the
region),
I don't think that can be avoided. The host doesn't have control over when the
CTAs (a gang) start -- they may even be serialized onto the same physical HW.
So each gang has to initialize its own instance. Or did you mean something else?
but more importantly, if the var is already mapped, you could
initialize the private var with old data.
Say
int arr[64];
// initialize arr
#pragma acc data copyin (arr)
{
// modify arr on the host
# pragma acc parallel firstprivate (arr)
{
...
}
}
Hm, I suspect that is either ill formed or the std does not contemplate.
Is that really what you want? If not, any reason not to implement
GOMP_MAP_FIRSTPRIVATE and GOMP_MAP_FIRSTPRIVATE_INT on the libgomp oacc-*
side and just use the OpenMP firstprivate handling in omp-low.c?
I would have to investigate ...
nathan
