public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
* [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable
@ 2021-04-13 8:58 burnus at gcc dot gnu.org
2021-04-13 13:16 ` [Bug middle-end/100059] " burnus at gcc dot gnu.org
` (6 more replies)
0 siblings, 7 replies; 8+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-04-13 8:58 UTC (permalink / raw)
To: gcc-bugs
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).
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
@ 2021-04-13 13:16 ` burnus at gcc dot gnu.org
2021-04-13 13:37 ` burnus at gcc dot gnu.org
` (5 subsequent siblings)
6 siblings, 0 replies; 8+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-04-13 13:16 UTC (permalink / raw)
To: gcc-bugs
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"
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
2021-04-13 13:16 ` [Bug middle-end/100059] " burnus at gcc dot gnu.org
@ 2021-04-13 13:37 ` burnus at gcc dot gnu.org
2021-04-13 14:11 ` burnus at gcc dot gnu.org
` (4 subsequent siblings)
6 siblings, 0 replies; 8+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-04-13 13:37 UTC (permalink / raw)
To: gcc-bugs
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"
};
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
2021-04-13 13:16 ` [Bug middle-end/100059] " burnus at gcc dot gnu.org
2021-04-13 13:37 ` burnus at gcc dot gnu.org
@ 2021-04-13 14:11 ` burnus at gcc dot gnu.org
2021-04-13 16:39 ` burnus at gcc dot gnu.org
` (3 subsequent siblings)
6 siblings, 0 replies; 8+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-04-13 14:11 UTC (permalink / raw)
To: gcc-bugs
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;
}
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
` (2 preceding siblings ...)
2021-04-13 14:11 ` burnus at gcc dot gnu.org
@ 2021-04-13 16:39 ` burnus at gcc dot gnu.org
2022-03-29 10:23 ` burnus at gcc dot gnu.org
` (2 subsequent siblings)
6 siblings, 0 replies; 8+ messages in thread
From: burnus at gcc dot gnu.org @ 2021-04-13 16:39 UTC (permalink / raw)
To: gcc-bugs
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
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
` (3 preceding siblings ...)
2021-04-13 16:39 ` burnus at gcc dot gnu.org
@ 2022-03-29 10:23 ` burnus at gcc dot gnu.org
2023-09-04 9:37 ` tschwinge at gcc dot gnu.org
2023-09-04 12:57 ` [Bug other/100059] " tschwinge at gcc dot gnu.org
6 siblings, 0 replies; 8+ messages in thread
From: burnus at gcc dot gnu.org @ 2022-03-29 10:23 UTC (permalink / raw)
To: gcc-bugs
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.
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug middle-end/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
` (4 preceding siblings ...)
2022-03-29 10:23 ` burnus at gcc dot gnu.org
@ 2023-09-04 9:37 ` tschwinge at gcc dot gnu.org
2023-09-04 12:57 ` [Bug other/100059] " tschwinge at gcc dot gnu.org
6 siblings, 0 replies; 8+ messages in thread
From: tschwinge at gcc dot gnu.org @ 2023-09-04 9:37 UTC (permalink / raw)
To: gcc-bugs
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.
^ permalink raw reply [flat|nested] 8+ messages in thread
* [Bug other/100059] [OpenMP] wrong code with 'declare target link' and a scalar variable
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
` (5 preceding siblings ...)
2023-09-04 9:37 ` tschwinge at gcc dot gnu.org
@ 2023-09-04 12:57 ` tschwinge at gcc dot gnu.org
6 siblings, 0 replies; 8+ messages in thread
From: tschwinge at gcc dot gnu.org @ 2023-09-04 12:57 UTC (permalink / raw)
To: gcc-bugs
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'"
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2023-09-04 12:57 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-04-13 8:58 [Bug middle-end/100059] New: [OpenMP] wrong code with 'declare target link' and a scalar variable burnus at gcc dot gnu.org
2021-04-13 13:16 ` [Bug middle-end/100059] " burnus at gcc dot gnu.org
2021-04-13 13:37 ` burnus at gcc dot gnu.org
2021-04-13 14:11 ` burnus at gcc dot gnu.org
2021-04-13 16:39 ` burnus at gcc dot gnu.org
2022-03-29 10:23 ` burnus at gcc dot gnu.org
2023-09-04 9:37 ` tschwinge at gcc dot gnu.org
2023-09-04 12:57 ` [Bug other/100059] " tschwinge at gcc dot gnu.org
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).