[PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Ilya Verbin]

Hi!

Currently if a target* pragma appears within a target region, GCC successfully
compiles such code (with a warning).  But the binary fails at run-time, since it
tries to call GOMP_target* functions on target.

The spec says: "If a target, target update, or target data construct appears
within a target region then the behavior is unspecified."

I see 2 options to make the behavior more user-friendly:
1. To return an error at compile-time.
2. To check at run-time in libgomp whether GOMP_target* is called on target, and
perform target-fallback if so.

If we will select option #1, the patch is ready.


diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index eaad52a..ae8b90a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2682,8 +2682,14 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 	}
       break;
     case GIMPLE_OMP_TARGET:
-      for (; ctx != NULL; ctx = ctx->outer)
-	if (is_targetreg_ctx (ctx))
+      {
+	bool target_p = false;
+	if (cgraph_node::get (current_function_decl)->offloadable)
+	  target_p = true;
+	for (; ctx != NULL && !target_p; ctx = ctx->outer)
+	  if (is_targetreg_ctx (ctx))
+	    target_p = true;
+	if (target_p)
 	  {
 	    const char *name;
 	    switch (gimple_omp_target_kind (stmt))
@@ -2693,9 +2699,11 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 	      case GF_OMP_TARGET_KIND_UPDATE: name = "target update"; break;
 	      default: gcc_unreachable ();
 	      }
-	    warning_at (gimple_location (stmt), 0,
-			"%s construct inside of target region", name);
+	    error_at (gimple_location (stmt),
+		      "%s construct inside of target region", name);
+	    return false;
 	  }
+      }
       break;
     default:
       break;
diff --git a/gcc/testsuite/c-c++-common/gomp/nesting-error-1.c b/gcc/testsuite/c-c++-common/gomp/nesting-error-1.c
new file mode 100644
index 0000000..ff6a75e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/nesting-error-1.c
@@ -0,0 +1,23 @@
+extern int i;
+
+void
+f_omp_target (void)
+{
+#pragma omp target
+  {
+#pragma omp target /* { dg-error "target construct inside of target region" } */
+    ;
+#pragma omp target data /* { dg-error "target data construct inside of target region" } */
+    ;
+#pragma omp target update to(i) /* { dg-error "target update construct inside of target region" } */
+
+#pragma omp parallel
+    {
+#pragma omp target /* { dg-error "target construct inside of target region" } */
+      ;
+#pragma omp target data /* { dg-error "target data construct inside of target region" } */
+      ;
+#pragma omp target update to(i) /* { dg-error "target update construct inside of target region" } */
+    }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c b/gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
deleted file mode 100644
index c39dd49..0000000
--- a/gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c
+++ /dev/null
@@ -1,23 +0,0 @@
-extern int i;
-
-void
-f_omp_target (void)
-{
-#pragma omp target
-  {
-#pragma omp target /* { dg-warning "target construct inside of target region" } */
-    ;
-#pragma omp target data /* { dg-warning "target data construct inside of target region" } */
-    ;
-#pragma omp target update to(i) /* { dg-warning "target update construct inside of target region" } */
-
-#pragma omp parallel
-    {
-#pragma omp target /* { dg-warning "target construct inside of target region" } */
-      ;
-#pragma omp target data /* { dg-warning "target data construct inside of target region" } */
-      ;
-#pragma omp target update to(i) /* { dg-warning "target update construct inside of target region" } */
-    }
-  }
-}
diff --git a/gcc/testsuite/gfortran.dg/gomp/target3.f90 b/gcc/testsuite/gfortran.dg/gomp/target3.f90
index 53a9682..e451548 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target3.f90
@@ -4,9 +4,18 @@
 subroutine foo (r)
   integer :: i, r
   !$omp target
-  !$omp target teams distribute parallel do reduction (+: r) ! { dg-warning "target construct inside of target region" }
+  !$omp target teams distribute parallel do reduction (+: r) ! { dg-error "target construct inside of target region" }
     do i = 1, 10
       r = r + 1
     end do
   !$omp end target
 end subroutine
+
+subroutine bar (r)
+  !$omp declare target
+  integer :: i, r
+  !$omp target teams distribute parallel do reduction (+: r) ! { dg-error "target construct inside of target region" }
+    do i = 1, 10
+      r = r + 1
+    end do
+end subroutine


  -- Ilya

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Ilya Verbin]

Hi Jakub!

We have 3 pending patches with warnings/errors about omp pragmas:

1. https://gcc.gnu.org/ml/gcc-patches/2015-01/msg00617.html
2. https://gcc.gnu.org/ml/gcc-patches/2015-01/msg00621.html
3. This one.

What should we do with them?

[ ] Rebase and continue pinging.
[ ] Postpone until Stage 1.
[ ] Completely remove.


On 12 Jan 00:22, Ilya Verbin wrote:
> Currently if a target* pragma appears within a target region, GCC successfully
> compiles such code (with a warning).  But the binary fails at run-time, since it
> tries to call GOMP_target* functions on target.
> 
> The spec says: "If a target, target update, or target data construct appears
> within a target region then the behavior is unspecified."
> 
> I see 2 options to make the behavior more user-friendly:
> 1. To return an error at compile-time.
> 2. To check at run-time in libgomp whether GOMP_target* is called on target, and
> perform target-fallback if so.
> 
> If we will select option #1, the patch is ready.

  -- Ilya

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Jakub Jelinek]

On Mon, Jan 12, 2015 at 12:22:44AM +0300, Ilya Verbin wrote:
> Currently if a target* pragma appears within a target region, GCC successfully
> compiles such code (with a warning).  But the binary fails at run-time, since it
> tries to call GOMP_target* functions on target.
> 
> The spec says: "If a target, target update, or target data construct appears
> within a target region then the behavior is unspecified."
> 
> I see 2 options to make the behavior more user-friendly:
> 1. To return an error at compile-time.
> 2. To check at run-time in libgomp whether GOMP_target* is called on target, and
> perform target-fallback if so.
> 
> If we will select option #1, the patch is ready.

Option #1 is just wrong.  There is nothing wrong with such constructs
appearing in #pragma omp declare target functions etc., the problem is
if you hit them at runtime.  You can very well have say #pragma omp declare
target function, that optionally invokes #pragma omp target region e.g.
based on its parameters, state of global variables, what other functions
return etc. - and the program can be written so that that condition just
never happens if the function is already offloaded.

	Jakub

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Ilya Verbin]

On 02 Feb 13:15, Jakub Jelinek wrote:
> On Mon, Jan 12, 2015 at 12:22:44AM +0300, Ilya Verbin wrote:
> > Currently if a target* pragma appears within a target region, GCC successfully
> > compiles such code (with a warning).  But the binary fails at run-time, since it
> > tries to call GOMP_target* functions on target.
> > 
> > The spec says: "If a target, target update, or target data construct appears
> > within a target region then the behavior is unspecified."
> > 
> > I see 2 options to make the behavior more user-friendly:
> > 1. To return an error at compile-time.
> > 2. To check at run-time in libgomp whether GOMP_target* is called on target, and
> > perform target-fallback if so.
> > 
> > If we will select option #1, the patch is ready.
> 
> Option #1 is just wrong.  There is nothing wrong with such constructs
> appearing in #pragma omp declare target functions etc., the problem is
> if you hit them at runtime.  You can very well have say #pragma omp declare
> target function, that optionally invokes #pragma omp target region e.g.
> based on its parameters, state of global variables, what other functions
> return etc. - and the program can be written so that that condition just
> never happens if the function is already offloaded.

I thought that "If a target, target update, or target data construct appears
within a target region then the behavior is unspecified." applies to '#pragma
omp declare target' functions as well, but evidently this applies only to
'#pragma omp target' regions.

But there is another issue, I forgot to mention it in the first mail.
Here is a testcase:

int main ()
{
  #pragma omp target
    {
      int x;
      #pragma omp target map(to: x)
	x;
    }
}

This causes an ICE in the offload compiler, since .omp_data_sizes.3 and
.omp_data_kinds.4 are used in main._omp_fn.0, which should be compiled for
target, but these variables are static without 'declare target' attribute.

main ()
{
  struct .omp_data_t.1 .omp_data_arr.2;
  static long unsigned int .omp_data_sizes.3[1] = {4};
  static unsigned char .omp_data_kinds.4[1] = {17};
  GOMP_target (-1, main._omp_fn.0, 0B, 0, 0B, 0B, 0B);
}

main._omp_fn.0 (void * .omp_data_i)
{
  struct .omp_data_t.1 .omp_data_arr.2;
  int x;
  .omp_data_arr.2.x = &x;
  GOMP_target (-1, main._omp_fn.1, 0B, 1, &.omp_data_arr.2,
	       &.omp_data_sizes.3, &.omp_data_kinds.4);
}

main._omp_fn.1 (struct .omp_data_t.1 & restrict .omp_data_i)
{
  int x [value-expr: *.omp_data_i->x];
}

Therefore I wanted just to forbid nested target regions.  Or should we make
omp_data_sizes and omp_data_kinds non-static?

  -- Ilya

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Martin Jambor]

Hi,

On Mon, Jan 12, 2015 at 12:22:44AM +0300, Ilya Verbin wrote:
> Hi!
> 
> Currently if a target* pragma appears within a target region, GCC successfully
> compiles such code (with a warning).  But the binary fails at run-time, since it
> tries to call GOMP_target* functions on target.
> 
> The spec says: "If a target, target update, or target data construct appears
> within a target region then the behavior is unspecified."
> 
> I see 2 options to make the behavior more user-friendly:
> 1. To return an error at compile-time.
> 2. To check at run-time in libgomp whether GOMP_target* is called on target, and
> perform target-fallback if so.
> 

What actually happens when an accelerator calls a libgomp function?
Is a target libgomp port invoked?  If so, it should easily know it
runs on a target even without a run-time check, I suppose.  Or do you
somehow bring that call back to the host?

Thanks,

Martin

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Jakub Jelinek]

On Tue, Feb 10, 2015 at 11:16:22AM +0100, Martin Jambor wrote:
> On Mon, Jan 12, 2015 at 12:22:44AM +0300, Ilya Verbin wrote:
> > Currently if a target* pragma appears within a target region, GCC successfully
> > compiles such code (with a warning).  But the binary fails at run-time, since it
> > tries to call GOMP_target* functions on target.
> > 
> > The spec says: "If a target, target update, or target data construct appears
> > within a target region then the behavior is unspecified."
> > 
> > I see 2 options to make the behavior more user-friendly:
> > 1. To return an error at compile-time.
> > 2. To check at run-time in libgomp whether GOMP_target* is called on target, and
> > perform target-fallback if so.
> > 
> 
> What actually happens when an accelerator calls a libgomp function?
> Is a target libgomp port invoked?  If so, it should easily know it
> runs on a target even without a run-time check, I suppose.  Or do you
> somehow bring that call back to the host?

The spec says that it is undefined behavior to invoke
#pragma omp target {,data,update} from within #pragma omp target region.
For intelmic, the offloading shared libraries are normally linked against
-lgomp and thus can call any functions from there.
For nvptx, libgomp still needs to be ported to that target.
So, what we can do is e.g. ignore the nested #pragma omp target* regions
inside of #pragma omp target, or turn them into __builtin_trap ().

	Jakub

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 2505 bytes --]

Hi!

On Tue, 10 Feb 2015 11:20:24 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Tue, Feb 10, 2015 at 11:16:22AM +0100, Martin Jambor wrote:
> > On Mon, Jan 12, 2015 at 12:22:44AM +0300, Ilya Verbin wrote:
> > > Currently if a target* pragma appears within a target region, GCC successfully
> > > compiles such code (with a warning).  But the binary fails at run-time, since it
> > > tries to call GOMP_target* functions on target.
> > > 
> > > The spec says: "If a target, target update, or target data construct appears
> > > within a target region then the behavior is unspecified."
> > > 
> > > I see 2 options to make the behavior more user-friendly:
> > > 1. To return an error at compile-time.
> > > 2. To check at run-time in libgomp whether GOMP_target* is called on target, and
> > > perform target-fallback if so.
> > > 
> > 
> > What actually happens when an accelerator calls a libgomp function?
> > Is a target libgomp port invoked?  If so, it should easily know it
> > runs on a target even without a run-time check, I suppose.  Or do you
> > somehow bring that call back to the host?
> 
> The spec says that it is undefined behavior to invoke
> #pragma omp target {,data,update} from within #pragma omp target region.

We're not currently implementing that, but let me mention that OpenACC
describes a concept of nested parallelism:

OpenACC 2.0a, 1.2 Execution Model:

    [...]
    On some devices, the accelerator may also create and launch parallel kernels, allowing for
    nested parallelism. In that case, the OpenACC directives may be executed by a host thread or
    an accelerator thread. [...]

OpenACC 2.0a, 2.6 Data Environment:

    [...] When a
    nested OpenACC construct is executed on the device, the default target device for that
    construct is the same device on which the encountering accelerator thread is executing. In
    that case, the target device shares memory with the encountering thread.

For PTX, this would use CUDA's Dynamic Parallelism,
<http://devblogs.nvidia.com/parallelforall/introduction-cuda-dynamic-parallelism/>,
for example.

> For intelmic, the offloading shared libraries are normally linked against
> -lgomp and thus can call any functions from there.
> For nvptx, libgomp still needs to be ported to that target.
> So, what we can do is e.g. ignore the nested #pragma omp target* regions
> inside of #pragma omp target, or turn them into __builtin_trap ().


Grüße,
 Thomas

[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]

Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions

[Ilya Verbin]

On Tue, Feb 10, 2015 at 11:20:24 +0100, Jakub Jelinek wrote:
> So, what we can do is e.g. ignore the nested #pragma omp target* regions
> inside of #pragma omp target, or turn them into __builtin_trap ().

I like this idea to ignore them at compile-time, but how to be with declare
target functions?

#pragma omp declare target
void foo (int x, int y)
{
  #pragma omp target map(to: y) if(x)
    y;
}
#pragma omp end declare target

This code seems to be correct, but it has the same problem with static
omp_data_sizes and omp_data_kinds.

  -- Ilya