diff mbox series

[OpenMP,5.0] Basic support of release/delete clauses on target/target-data directives

Message ID fc0a8875-b427-246c-c66e-b9d85e03a55d@codesourcery.com
State New
Headers show
Series [OpenMP,5.0] Basic support of release/delete clauses on target/target-data directives | expand

Commit Message

Chung-Lin Tang Dec. 16, 2020, 3:06 p.m. UTC 
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  <cltang@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.

Comments

Jakub Jelinek Jan. 11, 2021, 10:28 a.m. UTC | #1 
On Wed, Dec 16, 2020 at 11:06:10PM +0800, Chung-Lin Tang wrote:
> we have some other sollve_vv tests for OpenMP 5.0, which tests the occurrence of
> 'delete' map kind on the target directive.

That is strange.
In OpenMP 5.0, I certainly see the target data:
[163:16] A map-type in a map clause must be to, from, tofrom or alloc.
restriction and similarly for target:
[174:27] A map-type in a map clause must be to, from, tofrom or alloc.
restriction.
OpenMP 4.5 has those restrictions in [97:15] and [106:13].
So if sollve_vv tests use some other map-types in target or target data
construct map clauses, it should be reported to them and fixed there.

	Jakub
Chung-Lin Tang Jan. 12, 2021, 11:25 a.m. UTC | #2 
On 2021/1/11 6:28 PM, Jakub Jelinek wrote:
> On Wed, Dec 16, 2020 at 11:06:10PM +0800, Chung-Lin Tang wrote:
>> we have some other sollve_vv tests for OpenMP 5.0, which tests the occurrence of
>> 'delete' map kind on the target directive.
> 
> That is strange.
> In OpenMP 5.0, I certainly see the target data:
> [163:16] A map-type in a map clause must be to, from, tofrom or alloc.
> restriction and similarly for target:
> [174:27] A map-type in a map clause must be to, from, tofrom or alloc.
> restriction.

Thanks! I was originally trying to find such a restriction too, but missed those lines above.
Appears that this patch isn't needed then.

Chung-Lin
diff mbox series

Patch

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: