On 05/25/2016 09:33 AM, Thomas Schwinge wrote:
>> --- gcc/c/c-parser.c
>> +++ gcc/c/c-parser.c
>> @@ -13602,6 +13602,7 @@ c_parser_oacc_declare (c_parser *parser)
>>
>> switch (OMP_CLAUSE_MAP_KIND (t))
>> {
>> + case GOMP_MAP_FIRSTPRIVATE_POINTER:
>> case GOMP_MAP_FORCE_ALLOC:
>> case GOMP_MAP_FORCE_TO:
>> case GOMP_MAP_FORCE_DEVICEPTR:
> | case GOMP_MAP_DEVICE_RESIDENT:
> | break;
> |
> | case GOMP_MAP_POINTER:
> | /* Generated by c_finish_omp_clauses from array sections;
> | avoid spurious diagnostics. */
> | break;
>
> Is "case GOMP_MAP_FIRSTPRIVATE_POINTER" meant to replace the "case
> GOMP_MAP_POINTER"? If yes, then please remove that one (does that become
> gcc_unreachable?), and update/move the comment, or if not, please update
> the comment, too. ;-)
This can be pruned out. I'm preparing a follow up patch with it's removal.
>> --- gcc/c/c-typeck.c
>> +++ gcc/c/c-typeck.c
>
>> /* Handle array sections for clause C. */
>>
>> static bool
>> -handle_omp_array_sections (tree c, bool is_omp)
>> +handle_omp_array_sections (tree c, enum c_omp_region_type ort)
>> {
>> [...]
>> @@ -12427,7 +12427,7 @@ handle_omp_array_sections (tree c, bool is_omp)
>> && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
>> return false;
>> gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
>> - if (is_omp)
>> + if (ort == C_ORT_OMP || ort == C_ORT_ACC)
>> switch (OMP_CLAUSE_MAP_KIND (c))
>> {
>> case GOMP_MAP_ALLOC:
> | case GOMP_MAP_TO:
> | case GOMP_MAP_FROM:
> | case GOMP_MAP_TOFROM:
> | case GOMP_MAP_ALWAYS_TO:
> | case GOMP_MAP_ALWAYS_FROM:
> | case GOMP_MAP_ALWAYS_TOFROM:
> | case GOMP_MAP_RELEASE:
> | case GOMP_MAP_DELETE:
> | OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
> | break;
> | default:
>> break;
>> }
>
> Why doesn't that apply also to the other (OpenACC) map kinds? Comparing
> to the full list in include/gomp-constants.h:enum gomp_map_kind, there
> are several missing here.
It does look like there are situations where OpenACC can take
zero-length arrays, e.g. when the length argument is a variable. This
will be fixed in my follow up patch.
>> tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
>> - if (!is_omp)
>> + if (ort != C_ORT_OMP && ort != C_ORT_ACC)
>> OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
>> else if (TREE_CODE (t) == COMPONENT_REF)
>> OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
>
>> --- gcc/cp/parser.c
>> +++ gcc/cp/parser.c
>> @@ -35214,6 +35214,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token
>> *pragma_tok)
>> gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP);
>> switch (OMP_CLAUSE_MAP_KIND (t))
>> {
>> + case GOMP_MAP_FIRSTPRIVATE_POINTER:
>> case GOMP_MAP_FORCE_ALLOC:
>> case GOMP_MAP_FORCE_TO:
>> case GOMP_MAP_FORCE_DEVICEPTR:
>
> Likewise to my gcc/c/c-parser.c comments.
>
>> --- gcc/cp/semantics.c
>> +++ gcc/cp/semantics.c
>
>> /* Handle array sections for clause C. */
>>
>> static bool
>> -handle_omp_array_sections (tree c, bool is_omp)
>> +handle_omp_array_sections (tree c, enum c_omp_region_type ort)
>> {
>> [...]
>> @@ -4988,7 +4989,7 @@ handle_omp_array_sections (tree c, bool is_omp)
>> || (TREE_CODE (t) == COMPONENT_REF
>> && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
>> return false;
>> - if (is_omp)
>> + if (ort == C_ORT_OMP || ort == C_ORT_ACC)
>> switch (OMP_CLAUSE_MAP_KIND (c))
>> {
>> case GOMP_MAP_ALLOC:
>
> Likewise to my gcc/c/c-typeck.c comments.
>
>> @@ -5007,7 +5008,7 @@ handle_omp_array_sections (tree c, bool is_omp)
>> }
>> tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
>> OMP_CLAUSE_MAP);
>> - if (!is_omp)
>> + if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
>> OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
>> else if (TREE_CODE (t) == COMPONENT_REF)
>> OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
>
> Shouldn't that simply be "ort != C_ORT_OMP && ort != C_ORT_ACC"?
No, because then that wouldn't cover omp declare simd.
>> @@ -6054,7 +6070,7 @@ finish_omp_clauses (tree clauses, enum
>> c_omp_region_type ort)
>> omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
>> else
>> t = OMP_CLAUSE_DECL (c);
>> - if (t == current_class_ptr)
>> + if (ort != C_ORT_ACC && t == current_class_ptr)
>> {
>> error ("%<this%> allowed in OpenMP only in %<declare simd%>"
>> " clauses");
>
> ;-) Hmm, reminds me of the unresolved task to support the C++ "this"
> pointer in OpenACC... Anyway, in GCC trunk, we're not allowing "this"
> usage, I think, so I suppose this should stay as-is? (Possibly with an
> OpenACC-specific error message.)
What do you want to do about c++'s 'this' in OpenACC? It looks like
gomp4 partially supports it. Maybe we should wait to get clarification
from the technical committee?
>> @@ -6681,7 +6701,7 @@ finish_omp_clauses (tree clauses, enum
>> c_omp_region_type ort)
>> omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
>> remove = true;
>> }
>> - else if (t == current_class_ptr)
>> + else if (ort != C_ORT_ACC && t == current_class_ptr)
>> {
>> error ("%<this%> allowed in OpenMP only in %<declare simd%>"
>> " clauses");
>
> Likewise.
>
>> --- gcc/gimplify.c
>> +++ gcc/gimplify.c
>> @@ -6280,6 +6280,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx,
>> tree decl, bool in_code)
>> error ("variable %qE declared in enclosing "
>> "%<host_data%> region", DECL_NAME (decl));
>> nflags |= GOVD_MAP;
>> + if (octx->region_type == ORT_ACC_DATA
>> + && (n2->value & GOVD_MAP_0LEN_ARRAY))
>> + nflags |= GOVD_MAP_0LEN_ARRAY;
>> goto found_outer;
>> }
>> }
>
> Later on, everyone will have a hard time to understand that logic, so
> please add comments for such special handling. Why is ORT_ACC_DATA being
> handled differently from the OpenMP target data construct, for example?
It's because pointers in OpenACC have two separate meanings depending on
context, whereas in OpenMP they only have one meaning. In OpenACC, a
pointer by itself, e.g., copy(ptr), is supposed to be treated like
scalar. However, a pointer to a subarray, e.g. copy(ptr[0:10]) is
supposed to treated as firstprivate so that only the data gets the data
clause.
These gimplifier changes were necessary to get proper implicit data
clauses for parallel and kernels regions nested inside data regions. E.g.
#pragma acc data copy(ptr[0:10])
{
#pragma acc parallel
{
... ptr[x] = ...
}
}
Without these changes, ptr would be treated like a scalar.
>> @@ -6855,9 +6858,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
>> *pre_p,
>> {
>> case OMP_TARGET:
>> break;
>> + case OACC_DATA:
>> + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
>> + break;
>
> Likewise.
>
> Also add a "/* FALLTHRU */" comment here.
>
>> case OMP_TARGET_DATA:
>> case OMP_TARGET_ENTER_DATA:
>> case OMP_TARGET_EXIT_DATA:
>> + case OACC_ENTER_DATA:
>> + case OACC_EXIT_DATA:
>> case OACC_HOST_DATA:
>> if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
>> || (OMP_CLAUSE_MAP_KIND (c)
> | == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
> | /* For target {,enter ,exit }data only the array slice is
> | mapped, but not the pointer to it. */
> | remove = true;
> | break;
> | default:
> | break;
> | }
>
> By the way, why is this not relevant for the OpenACC update and OpenMP
> target update directives, OACC_UPDATE and OMP_TARGET_UPDATE? Is it
> because theses only update existing mappings but don't create new ones?
I suppose they can be added here, but lower_omp_target already ignores
GOMP_MAP_FIRSTPRIVATE_POINTER for non-offloaded and data_region regions.
Cesar
