On 21 Nov 21:36, Jakub Jelinek wrote:
> On Fri, Nov 21, 2014 at 11:19:26PM +0300, Ilya Verbin wrote:
> > '#pragma omp critical (name)' can be placed in the function, marked
> > with '#pragma omp declare target', in this case the corresponding node
> > should be marked as offloadable too.
> > Bootstrapped/regtested on x86_64-linux and i686-linux, ok for trunk?
>
> Please add a testcase for this.
Here is the testcase, based on critical-2.c, ok for trunk?
gcc/
* omp-low.c (lower_omp_critical): Mark critical sections
inside target functions as offloadable.
libgomp/
* testsuite/libgomp.c/target-critical-1.c: New test.
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index dff1504..621958c 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -9387,16 +9387,6 @@ lower_omp_critical (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
DECL_ARTIFICIAL (decl) = 1;
DECL_IGNORED_P (decl) = 1;
- /* If '#pragma omp critical' is inside target region, the symbol must
- be marked for offloading. */
- omp_context *octx;
- for (octx = ctx->outer; octx; octx = octx->outer)
- if (is_targetreg_ctx (octx))
- {
- varpool_node::get_create (decl)->offloadable = 1;
- break;
- }
-
varpool_node::finalize_decl (decl);
critical_name_mutexes->put (name, decl);
@@ -9404,6 +9394,20 @@ lower_omp_critical (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
else
decl = *n;
+ /* If '#pragma omp critical' is inside target region or
+ inside function marked as offloadable, the symbol must be
+ marked as offloadable too. */
+ omp_context *octx;
+ if (cgraph_node::get (current_function_decl)->offloadable)
+ varpool_node::get_create (decl)->offloadable = 1;
+ else
+ for (octx = ctx->outer; octx; octx = octx->outer)
+ if (is_targetreg_ctx (octx))
+ {
+ varpool_node::get_create (decl)->offloadable = 1;
+ break;
+ }
+
lock = builtin_decl_explicit (BUILT_IN_GOMP_CRITICAL_NAME_START);
lock = build_call_expr_loc (loc, lock, 1, build_fold_addr_expr_loc (loc,
decl));
diff --git a/libgomp/testsuite/libgomp.c/target-critical-1.c
b/libgomp/testsuite/libgomp.c/target-critical-1.c
new file mode 100644
index 0000000..84ad558
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-critical-1.c
@@ -0,0 +1,72 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+#define N 2000
+
+#pragma omp declare target
+int foo ()
+{
+ int A[N];
+ int i, nthreads;
+ int res = 0;
+
+ #pragma omp parallel shared (A, nthreads)
+ {
+ #pragma omp master
+ nthreads = omp_get_num_threads ();
+
+ #pragma omp for
+ for (i = 0; i < N; i++)
+ A[i] = 0;
+
+ #pragma omp critical (crit1)
+ for (i = 0; i < N; i++)
+ A[i]++;
+ }
+
+ for (i = 0; i < N; i++)
+ if (A[i] != nthreads)
+ res = 1;
+
+ return res;
+}
+#pragma omp end declare target
+
+int main ()
+{
+ int res1, res2;
+
+ #pragma omp target map (from: res1, res2)
+ {
+ int B[N];
+ int i, nthreads;
+
+ res1 = foo ();
+
+ #pragma omp parallel shared (B, nthreads)
+ {
+ #pragma omp master
+ nthreads = omp_get_num_threads ();
+
+ #pragma omp for
+ for (i = 0; i < N; i++)
+ B[i] = 0;
+
+ #pragma omp critical (crit2)
+ for (i = 0; i < N; i++)
+ B[i]++;
+ }
+
+ res2 = 0;
+ for (i = 0; i < N; i++)
+ if (B[i] != nthreads)
+ res2 = 1;
+ }
+
+ if (res1 || res2)
+ abort ();
+
+ return 0;
+}
-- Ilya
