Add a restriction on allocate clause (OpenMP 5.0)
Commit Message
An allocate clause in target region must specify an allocator
unless the compilation unit has requires construct with
dynamic_allocators clause. Current implementation of the allocate
clause did not check for this restriction. This patch fills that
gap.
gcc/ChangeLog:
* omp-low.cc (omp_maybe_offloaded_ctx): New prototype.
(scan_sharing_clauses): Check a restriction on allocate clause.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/allocate-2.c: Add tests.
* c-c++-common/gomp/allocate-8.c: New test.
* gfortran.dg/gomp/allocate-3.f90: Add tests.
* gcc.dg/gomp/pr104517.c: Update.
---
gcc/omp-low.cc | 10 ++++++++++
gcc/testsuite/c-c++-common/gomp/allocate-2.c | 15 +++++++++++++++
gcc/testsuite/c-c++-common/gomp/allocate-8.c | 18 ++++++++++++++++++
gcc/testsuite/gcc.dg/gomp/pr104517.c | 18 ++++++++++--------
gcc/testsuite/gfortran.dg/gomp/allocate-3.f90 | 14 ++++++++++++++
5 files changed, 67 insertions(+), 8 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/allocate-8.c
Comments
On Fri, Feb 18, 2022 at 11:13:16PM +0000, Hafiz Abid Qadeer wrote:
> An allocate clause in target region must specify an allocator
> unless the compilation unit has requires construct with
> dynamic_allocators clause. Current implementation of the allocate
> clause did not check for this restriction. This patch fills that
> gap.
>
> gcc/ChangeLog:
>
> * omp-low.cc (omp_maybe_offloaded_ctx): New prototype.
> (scan_sharing_clauses): Check a restriction on allocate clause.
>
> gcc/testsuite/ChangeLog:
>
> * c-c++-common/gomp/allocate-2.c: Add tests.
> * c-c++-common/gomp/allocate-8.c: New test.
> * gfortran.dg/gomp/allocate-3.f90: Add tests.
> * gcc.dg/gomp/pr104517.c: Update.
I think it has even stronger requirement, that the allocator is present
and is present in uses_allocators clause.
But we don't parse uses_allocators, so this is a good step forward.
Ok.
Jakub
@@ -195,6 +195,7 @@ static vec<gomp_task *> task_cpyfns;
static void scan_omp (gimple_seq *, omp_context *);
static tree scan_omp_1_op (tree *, int *, void *);
+static bool omp_maybe_offloaded_ctx (omp_context *ctx);
#define WALK_SUBSTMTS \
case GIMPLE_BIND: \
@@ -1169,6 +1170,15 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
|| !integer_onep (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))
|| OMP_CLAUSE_ALLOCATE_ALIGN (c) != NULL_TREE))
{
+ /* The allocate clauses that appear on a target construct or on
+ constructs in a target region must specify an allocator expression
+ unless a requires directive with the dynamic_allocators clause
+ is present in the same compilation unit. */
+ if (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) == NULL_TREE
+ && ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
+ && omp_maybe_offloaded_ctx (ctx))
+ error_at (OMP_CLAUSE_LOCATION (c), "%<allocate%> clause must"
+ " specify an allocator here");
if (ctx->allocate_map == NULL)
ctx->allocate_map = new hash_map<tree, tree>;
tree val = integer_zero_node;
@@ -43,3 +43,18 @@ foo (int x, int z)
#pragma omp parallel private (x) allocate (0 : x) /* { dg-error "'allocate' clause allocator expression has type 'int' rather than 'omp_allocator_handle_t'" } */
bar (x, &x, 0);
}
+
+void
+foo1 ()
+{
+ int a = 10;
+#pragma omp target
+ {
+ #pragma omp parallel private (a) allocate(a) // { dg-error "'allocate' clause must specify an allocator here" }
+ a = 20;
+ }
+#pragma omp target private(a) allocate(a) // { dg-error "'allocate' clause must specify an allocator here" }
+ {
+ a = 30;
+ }
+}
new file mode 100644
@@ -0,0 +1,18 @@
+#pragma omp requires dynamic_allocators
+void
+foo ()
+{
+ int a = 10;
+#pragma omp parallel private (a) allocate(a)
+ a = 20;
+#pragma omp target
+ {
+ #pragma omp parallel private (a) allocate(a)
+ a = 30;
+ }
+#pragma omp target private(a) allocate(a)
+ {
+ a = 40;
+ }
+}
+
@@ -2,11 +2,13 @@
/* { dg-do compile } */
/* { dg-options "-O1 -fcompare-debug -fopenmp -fno-tree-ter -save-temps" } */
-enum {
- omp_default_mem_alloc,
- omp_large_cap_mem_alloc,
- omp_const_mem_alloc,
- omp_high_bw_mem_alloc
+typedef enum omp_allocator_handle_t
+{
+ omp_null_allocator = 0,
+ omp_default_mem_alloc = 1,
+ omp_large_cap_mem_alloc = 2,
+ omp_const_mem_alloc = 3,
+ omp_high_bw_mem_alloc = 4,
} omp_allocator_handle_t;
int t, bar_nte, bar_tl, bar_i3, bar_dd;
@@ -23,7 +25,7 @@ bar (int *idp, int s, int nth, int g, int nta, int fi, int pp, int *q,
int p = 0, i2 = 0, i1 = 0, m = 0, d = 0;
#pragma omp target parallel for \
- device(p) firstprivate (f) allocate (f) in_reduction(+:r2)
+ device(p) firstprivate (f) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
for (int i = 0; i < 4; i++)
ll++;
@@ -31,8 +33,8 @@ bar (int *idp, int s, int nth, int g, int nta, int fi, int pp, int *q,
device(d) map (m) \
if (target: p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
if (parallel: i2) reduction(+:r) num_threads (nth) linear (ll) \
- schedule(static) collapse(1) nowait depend(inout: d) allocate (f) \
- in_reduction(+:r2)
+ schedule(static) collapse(1) nowait depend(inout: d) \
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
for (int i = 0; i < 4; i++)
ll++;
@@ -12,3 +12,17 @@ subroutine foo(x)
!$omp end parallel do simd
end subroutine
+
+subroutine bar(a)
+ implicit none
+ integer :: a
+!$omp target
+ !$omp parallel private (a) allocate(a) ! { dg-error "'allocate' clause must specify an allocator here" }
+ a = 20
+ !$omp end parallel
+!$omp end target
+
+!$omp target private(a) allocate(a) ! { dg-error "'allocate' clause must specify an allocator here" }
+ a = 30;
+!$omp end target
+end subroutine