[Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable

[Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

            Bug ID: 100059
           Summary: [OpenMP] wrong code with 'declare target link' and a
                    scalar variable
           Product: gcc
           Version: 11.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---
            Target: nvptx-none

Created attachment 50578
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=50578&action=edit
Testcase, compile with 'gcc -fopenmp' with an (nvptx) offloading enabled
compiler

The attached testcase has four global vars:
  int a[N], b[N], c[N];
  int i = 0;

which are used with:

  #pragma omp declare target link(a,c,b,i)  

and

  #pragma omp target map(to: i) map(tofrom: a, b, c) 
    update();

After the call, the one array is not updated but has the previous host value:
* For 'declare target link(a,c,b,i)': not updated array is 'b'
* For 'declare target link(a,b,c,i)': not updated array is 'c'

Namely, the testcase shows:
i=5: A=5, B=6, C=7    // Original host value
i=5: A=6, B=6, C=10   // target call to 'update': 'B' is not updated
i=5: A=7, B=8, C=13   // host update, OK


Additionally, when using 'for (int i' in 'update': The error only shows up if
there is 'i'
And it occurs for the item before 'i', namely:
  link(a,c,b,i) → 'b' is wrong
  link(a,c,i,b) → 'c' is wrong


The issue occurs with nvptx. It works without offloading and I think it works
with GCN (did not show up in the test logs, but I have not manually verified).

[Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |vries at gcc dot gnu.org

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
For
  #pragma omp declare target link(a,c,b,i)
the vars are walked in reverse order (= in normal order of 'offload_vars'):

* In a.xnvptx-none.mkoffload.s:
------------
//:FUNC_MAP "main$_omp_fn$0"
//:VAR_MAP "i$linkptr"


// BEGIN GLOBAL VAR DEF: i$linkptr
        .visible .global .align 8 .u64 i$linkptr[1];
//:VAR_MAP "b$linkptr"

// BEGIN GLOBAL VAR DEF: b$linkptr
        .visible .global .align 8 .u64 b$linkptr[1];
//:VAR_MAP "c$linkptr"

// BEGIN GLOBAL VAR DEF: c$linkptr
        .visible .global .align 8 .u64 c$linkptr[1];
//:VAR_MAP "a$linkptr"

// BEGIN GLOBAL VAR DEF: a$linkptr
        .visible .global .align 8 .u64 a$linkptr[1];

// BEGIN GLOBAL FUNCTION DECL: gomp_nvptx_main
------------

So far so good. However, after:
  nvptx-none/bin/as -m sm_35 -o ./a.xnvptx-none.mkoffload.o
./a.xnvptx-none.mkoffload.s

the 'a.xnvptx-none.mkoffload.o' contains:

----------------
// BEGIN GLOBAL VAR DECL: __nvptx_uni
.extern .shared .u32 __nvptx_uni[32];
//:VAR_MAP "b$linkptr"
// BEGIN GLOBAL VAR DEF: b$linkptr
.visible .global .align 8 .u64 b$linkptr[1];
//:FUNC_MAP "main$_omp_fn$0"
//:VAR_MAP "i$linkptr"
// BEGIN GLOBAL VAR DEF: i$linkptr
.visible .global .align 8 .u64 i$linkptr[1];
//:VAR_MAP "c$linkptr"
// BEGIN GLOBAL VAR DEF: c$linkptr
.visible .global .align 8 .u64 c$linkptr[1];
//:VAR_MAP "a$linkptr"
// BEGIN GLOBAL VAR DEF: a$linkptr
.visible .global .align 8 .u64 a$linkptr[1];
// BEGIN GLOBAL FUNCTION DEF: update
----------------

Namely: 'b' and 'i' swapped the order!

The problem seems to be the misplaced:
  //:FUNC_MAP "main$_omp_fn$0"

[Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Probably a read herring as I still get for a.xnvptx-none.c:

static const char *const var_mappings[] = {
        "b$linkptr",
        "i$linkptr",
        "c$linkptr",
        "a$linkptr"
};

[Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
With GCN it does work:
 i=5: A=5, B=6, C=7
 i=5: A=6, B=8, C=10
 i=5: A=7, B=10, C=13

And the follow testcase prints:
* with nvptx:
  42, 5, 0, 0
* with amdgcn:
  5, 42, 43, 44

This something must be reversed ...
//-----------------------------------------------
int i = 0, a[3];

void update () { 
  i = 5; a[0] = 42; a[1] = 43; a[2] = 44;
}

#pragma omp declare target link(i,a)
#pragma omp declare target to(update)  

int main () {
  #pragma omp target map(from: i, a)
    update();
__builtin_printf("%d, %d, %d, %d\n", i,a[0],a[1],a[2]);
  return 0;
}

[Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Last reconfirmed|                            |2021-04-13
           Assignee|unassigned at gcc dot gnu.org      |burnus at gcc dot gnu.org
             Status|UNCONFIRMED                 |ASSIGNED
     Ever confirmed|0                           |1

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
That's in issue with how the nvptx pseudoassembler re-outputs the assembler,
which reorders the symbols – such that the order in var_mappings is wrong (i.e.
differs from the order on the host).

(Order is based on traversing the htab – I wonder why this did not show up
before?)  

See ISSUE DESCRIPTION and SOLUTION (PATCH) at
https://github.com/MentorEmbedded/nvptx-tools/pull/29

[Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

--- Comment #5 from Tobias Burnus <burnus at gcc dot gnu.org> ---
Another example for this is https://github.com/clang-ykt/omptests 's
t-same-name-definitions which has the same issue + is fixed by the pull request
for nvptx-tools.

[Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[tschwinge at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |tschwinge at gcc dot gnu.org
           See Also|                            |https://github.com/MentorEm
                   |                            |bedded/nvptx-tools/pull/29

--- Comment #6 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #4)
> https://github.com/MentorEmbedded/nvptx-tools/pull/29

This is now finally incorporated, sorry for the (long...) delay.

Are you going to push the GCC-level test case (submitted in this PR), or want
me to?  For nvptx offloading, it'll FAIL its execution test until nvptx-tools
updated, but that's OK in my opinion.

[Bug other/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable

[tschwinge at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100059

Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
          Component|middle-end                  |other
         Resolution|---                         |FIXED
             Status|ASSIGNED                    |RESOLVED
           Assignee|burnus at gcc dot gnu.org          |tschwinge at gcc dot gnu.org

--- Comment #7 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
commit r14-3650-gfe0f9e09413047484441468b05288412760d8a09 "Add
'libgomp.c-c++-common/pr100059-1.c'"