Created attachment 50578 [details] 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).
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'"