From 1a8c4d9ed36556a95bd7d53c04d2ec4c95594061 Mon Sep 17 00:00:00 2001 From: Hafiz Abid Qadeer Date: Fri, 18 Feb 2022 21:28:08 +0000 Subject: Add a restriction on allocate clause (OpenMP 5.0) 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 diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 4c52886..49481b2 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -195,6 +195,7 @@ static vec 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: \ @@ -1154,6 +1155,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), "% clause must" + " specify an allocator here"); if (ctx->allocate_map == NULL) ctx->allocate_map = new hash_map; tree val = integer_zero_node; diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-2.c b/gcc/testsuite/c-c++-common/gomp/allocate-2.c index cc77efc..6bb4a8a 100644 --- a/gcc/testsuite/c-c++-common/gomp/allocate-2.c +++ b/gcc/testsuite/c-c++-common/gomp/allocate-2.c @@ -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; + } +} diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-8.c b/gcc/testsuite/c-c++-common/gomp/allocate-8.c new file mode 100644 index 0000000..bacefaf --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/allocate-8.c @@ -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; + } +} + diff --git a/gcc/testsuite/gcc.dg/gomp/pr104517.c b/gcc/testsuite/gcc.dg/gomp/pr104517.c index efb3175..7e3bd1a 100644 --- a/gcc/testsuite/gcc.dg/gomp/pr104517.c +++ b/gcc/testsuite/gcc.dg/gomp/pr104517.c @@ -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++; diff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-3.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-3.f90 index 7b57be9..0bee99d 100644 --- a/gcc/testsuite/gfortran.dg/gomp/allocate-3.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/allocate-3.f90 @@ -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 -- cgit v1.1