Hi Jakub,
we have some other sollve_vv tests for OpenMP 5.0, which tests the occurrence of
'delete' map kind on the target directive.
Looking through the specification, as early as 4.5, it seems that there's no
wording that release/delete map kinds are explicitly prohibited on structured
block
device directives like target and target data (though may not be intuitive if
used in those cases)
According to the description of map clause operation, 'release' seems to be
essentially
the equivalent of 'alloc' if used this way.
The 'delete' map kind might be of some special use to "deep de-allocate" some
mappings on the way out of a target[data] region. Seems to be sometimes useful,
I think.
The attached patch is basically just a preliminary patch that lifts the
rejection of
release/delete map clauses on target/target-data; it no longer rejects those
kinds,
but has not yet any special handling of 'delete' in
gomp_map/unmap_vars_internal yet,
making them essentially just the same as 'alloc'. That said, it helps with some
testcase
compilation that doesn't test all the limitations.
(proper implementation of these clauses inside libgomp will be finished later)
Is this okay for trunk after stage1 reopens?
Thanks,
Chung-Lin
2020-12-16 Chung-Lin Tang <clt...@codesourcery.com>
gcc/c/
* c-parser.c (c_parser_omp_target_data): Allow 'release' and 'delete'
map kinds, update error message string.
(c_parser_omp_target): Likewise.
gcc/cp/
* parser.c (cp_parser_omp_target_data): Allow 'release' and 'delete'
map kinds, update error message string.
(cp_parser_omp_target): Likewise.
libgomp/
* target.c (gomp_map_vars_internal): Add GOMP_MAP_RELEASE and
GOMP_MAP_DELETE cases.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 69ecdb5e822..970e620b63c 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19563,26 +19563,28 @@ c_parser_omp_target_data (location_t loc, c_parser
*parser, bool *if_p)
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target data%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
+ "than %<to%>, %<from%>, %<tofrom%>, %<alloc%>, "
+ "%<release%>, or %<delete%> on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
continue;
}
else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
map_seen = 3;
pc = &OMP_CLAUSE_CHAIN (*pc);
@@ -20027,23 +20029,25 @@ check_clauses:
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
+ "than %<to%>, %<from%>, %<tofrom%>, %<alloc%>, "
+ "%<release%>, or %<delete%> on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
continue;
}
pc = &OMP_CLAUSE_CHAIN (*pc);
}
cfun->has_omp_target = true;
return true;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7ea8c28830e..55f3ab3f852 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -41430,27 +41430,29 @@ cp_parser_omp_target_data (cp_parser *parser,
cp_token *pragma_tok, bool *if_p)
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target data%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
+ "than %<to%>, %<from%>, %<tofrom%>, %<alloc%>, "
+ "%<release%>, or %<delete%> on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
continue;
}
else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
map_seen = 3;
pc = &OMP_CLAUSE_CHAIN (*pc);
@@ -41897,24 +41899,26 @@ check_clauses:
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
+ "than %<to%>, %<from%>, %<tofrom%>, %<alloc%>, "
+ "%<release%>, or %<delete%> on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
continue;
}
pc = &OMP_CLAUSE_CHAIN (*pc);
}
return true;
}
diff --git a/libgomp/target.c b/libgomp/target.c
index 6152f58e13d..c7d1b55bd6e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1280,14 +1280,16 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
switch (kind & typemask)
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
break;
case GOMP_MAP_TO:
case GOMP_MAP_TOFROM:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_TOFROM:
