Ilya Verbin - Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions (original) (raw)

This is the mail archive of the gcc-patches@gcc.gnu.orgmailing list for the GCC project.

• From: Ilya Verbin
• To: Jakub Jelinek
• Cc: gcc-patches at gcc dot gnu dot org, Kirill Yukhin , Thomas Schwinge
• Date: Mon, 9 Feb 2015 22:46:08 +0300
• Subject: Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions
• Authentication-results: sourceware.org; auth=none
• References: <20150111212244 dot GD30445 at msticlxl57 dot ims dot intel dot com> <20150202121534 dot GP1746 at tucnak dot redhat dot com>

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

• References:
  • Re: [PATCH][RFC][OpenMP] Forbid target* pragmas in target regions
    * From: Jakub Jelinek