Bug 100059

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"

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"
};

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;
}

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

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.

(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.

commit r14-3650-gfe0f9e09413047484441468b05288412760d8a09 "Add 'libgomp.c-c++-common/pr100059-1.c'"