diff options
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/omp-expand.cc | 14 | ||||
| -rw-r--r-- | gcc/omp-oacc-kernels-decompose.cc | 15 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/self-clause-2.c | 6 | ||||
| -rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 | 2 |
4 files changed, 27 insertions, 10 deletions
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index 8576b938102..5c6a7f2e381 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -10334,9 +10334,19 @@ expand_omp_target (struct omp_region *region) if ((c = omp_find_clause (clauses, OMP_CLAUSE_SELF)) != NULL_TREE) { - gcc_assert (is_gimple_omp_oacc (entry_stmt) && offloaded); + gcc_assert ((is_gimple_omp_oacc (entry_stmt) && offloaded) + || (gimple_omp_target_kind (entry_stmt) + == GF_OMP_TARGET_KIND_OACC_DATA_KERNELS)); - edge e = split_block_after_labels (new_bb); + edge e; + if (offloaded) + e = split_block_after_labels (new_bb); + else + { + gsi = gsi_last_nondebug_bb (new_bb); + gsi_prev (&gsi); + e = split_block (new_bb, gsi_stmt (gsi)); + } basic_block cond_bb = e->src; new_bb = e->dest; remove_edge (e); diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index ffc0a8f813e..dfbb34935d0 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -1519,17 +1519,18 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt) break; } } - else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF) + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF) { - /* If there is an 'if' clause, it must be duplicated to the - enclosing data region. Temporarily remove the if clause's - chain to avoid copying it. */ + /* If there is an 'if' or 'self' clause, it must be duplicated to the + enclosing data region. Temporarily remove its chain to avoid + copying it. */ tree saved_chain = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = NULL; - tree new_if_clause = unshare_expr (c); + tree new_clause = unshare_expr (c); OMP_CLAUSE_CHAIN (c) = saved_chain; - OMP_CLAUSE_CHAIN (new_if_clause) = data_clauses; - data_clauses = new_if_clause; + OMP_CLAUSE_CHAIN (new_clause) = data_clauses; + data_clauses = new_clause; } } /* Restore the original order of the clauses. */ diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c index 769694baec9..3ac29a03bc4 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c @@ -1,6 +1,8 @@ /* See also 'if-clause-2.c'. */ /* { dg-additional-options "-fdump-tree-gimple" } */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } + { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */ void f (short c) @@ -11,6 +13,8 @@ f (short c) #pragma acc kernels self(c) copy(c) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } + { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single async\(-1\) num_gangs\(1\) map\(force_present:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } */ ++c; #pragma acc serial self(c) copy(c) @@ -29,6 +33,8 @@ g (short d) #pragma acc kernels self copy(d) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "omp_oacc_kernels_decompose" } } + { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single async\(-1\) num_gangs\(1\) map\(force_present:d \[len: [0-9]+\]\) self\(1+\)$} 1 "omp_oacc_kernels_decompose" } } */ ++d; #pragma acc serial self copy(d) diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index 1ba04a84e12..2ee578f7f32 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -42,5 +42,5 @@ end program test ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } +! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\) self\(1\)$} 1 "omp_oacc_kernels_decompose" } } ! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } |