Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Ensure :VAR_MAP and :FUNC_MAP are output in order #29

Merged
merged 1 commit into from
Sep 4, 2023

Conversation

tob2
Copy link
Contributor

@tob2 tob2 commented Apr 13, 2021

Before, only if there is an '.extern', is_decl is set and only
then the variable/function is added to 'decls'. However, global
variables/functions may not have .extern - even if they show up
with :VAR_MAP and :FUNC_MAP. GCC's mkoffload uses the latter and
requires that the order matches the original order.

As the output is done as:

	write_stmts (out, rev_stmts (decls));
	htab_traverse (symbol_table, traverse, (void *)out);
	write_stmts (out, rev_stmts (fns));

it might have happen that :VAR_MAP/:FUNC_MAP lines will be output
in hash-map order, which can cause the wrong order.

https://gcc.gnu.org/PR100059

  • nvptx-as.c (parse_file): Set is_decl when there is a :VAR_MAP or :FUNC_MAP.

Before, only if there is an '.extern', is_decl is set and only
then the variable/function is added to 'decls'. However, global
variables/functions may not have .extern - even if they show up
with :VAR_MAP and :FUNC_MAP. GCC's mkoffload uses the latter and
requires that the order matches the original order.

As the output is done as:
  write_stmts (out, rev_stmts (decls));
  htab_traverse (symbol_table, traverse, (void *)out);
  write_stmts (out, rev_stmts (fns));
it might have happen that :VAR_MAP/:FUNC_MAP lines will be output
in hash-map order, which can cause the wrong order.

	https://gcc.gnu.org/PR100059
	* nvptx-as.c (parse_file): Set is_decl when there is a
	:VAR_MAP or :FUNC_MAP.
@tob2 tob2 requested a review from tschwinge April 13, 2021 16:28
@tob2
Copy link
Contributor Author

tob2 commented Apr 13, 2021

Testcase is:

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;
}
  • With the patch (or host compilation or GCN compilation), we get: 5, 42, 43, 44
  • Without the patch, we get: 42, 5, 0, 0

By inserting comment outputs before the write_stmts/htab_traverse it shows that both :VAR_MAP and :FUNC_MAP are output during the traverse!

Once this patch is in, the longer example at https://gcc.gnu.org/bugzilla/attachment.cgi?id=50578 (https://gcc.gnu.org/PR100059) should be committed to GCC.

Side question: Why did it not show up before?

@vries
Copy link
Contributor

vries commented Apr 22, 2021

Testcase is:

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;
}
* **With the patch** (or host compilation or GCN compilation), we get: `5, 42, 43, 44`

* **Without the patch**, we get: `42, 5, 0, 0`

Reproduced both the without and with patch behaviour.

@vries
Copy link
Contributor

vries commented Mar 2, 2022

Thomas, ping.

@tob2
Copy link
Contributor Author

tob2 commented Mar 29, 2022

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.

@tob2
Copy link
Contributor Author

tob2 commented Apr 11, 2022

@tschwinge — Ping²

@vries
Copy link
Contributor

vries commented Dec 14, 2022

@tschwinge — Ping^3

@tschwinge tschwinge merged commit 1b5946d into SourceryTools:master Sep 4, 2023
tschwinge added a commit that referenced this pull request Sep 6, 2023
#29]

... documenting (differently) buggy 'as' behavior re
commit 1b5946d,
commit 26095fd
"Ensure :VAR_MAP and :FUNC_MAP are output in order": GCC/nvptx offloading via
x86_64-linux-gnu vs. powerpc64le-linux-gnu host emit slightly different nvptx
code, which, for example, in 'test/as/order-2-x86_64-linux-gnu.o' vs.
'test/as/order-2-powerpc64le-linux-gnu.o', now loses an initializer:

    @@ -8,4 +8,6 @@
     //:FUNC_MAP "main$_omp_fn$0"
     //:VAR_MAP "data"
    +// BEGIN GLOBAL VAR DEF: data
    +.visible .global .align 4 .u32 data[1] =
     // BEGIN GLOBAL FUNCTION DECL: gomp_nvptx_main
     .extern .func gomp_nvptx_main (.param .u64 %in_ar1, .param .u64 %in_ar2);
    @@ -14,7 +16,4 @@
     // BEGIN GLOBAL VAR DECL: __nvptx_uni
     .extern .shared .u32 __nvptx_uni[32];
    -// BEGIN GLOBAL VAR DEF: data
    -.visible .global .align 4 .u32 data[1] =
    -{5 };
     // END PREAMBLE
     .visible .entry main$_omp_fn$0 (.param .u64 %arg, .param .u64 %stack, .param .u64 %sz)

..., and thus:

    [...]
    ptxas /tmp/ccFhz3dP.o, line 12; fatal   : Parsing error near '.extern': syntax error
    ptxas fatal   : Ptx assembly aborted due to errors
    nvptx-as: ptxas returned 255 exit status
    nvptx mkoffload: fatal error: [...]/powerpc64le-unknown-linux-gnu-accel-nvptx-none-gcc returned 1 exit status
    [...]
@tschwinge
Copy link
Member

Note commit 3f67564 "Ensure :VAR_MAP and :FUNC_MAP are output in order: add more test cases [#29]".

tschwinge added a commit that referenced this pull request Sep 6, 2023
Fix the 'as' issue described in commit 3f67564
"Ensure :VAR_MAP and :FUNC_MAP are output in order: add more test cases [#29]"
by handling '//:VAR_MAP', '//:FUNC_MAP' as stand-alone directives
instead of 'V_comment | V_prefix_comment'.

This also happens to unify the x86_64-linux-gnu, powerpc64le-linux-gnu
'test/as/order-1-{x86_64-linux-gnu,powerpc64le-linux-gnu}.o' as well as
'test/as/order-2-{x86_64-linux-gnu,powerpc64le-linux-gnu}.o' files.
@tschwinge
Copy link
Member

Note commit 3f67564 "Ensure :VAR_MAP and :FUNC_MAP are output in order: add more test cases [#29]".

..., which is now resolved via commit aa3404a "Ensure :VAR_MAP and :FUNC_MAP are output in order, part II [#29]".

tschwinge pushed a commit that referenced this pull request Apr 19, 2024
Similar to #29 for `:VAR_MAP` and `:FUNC_MAP` (cf. https://gcc.gnu.org/PR100059),
nvptx-as must keep the order the same for the entries of newly added indirect-function map table, for which the new tag `:IND_FUNC_MAP` is used.

**Cross ref:** https://gcc.gnu.org/pipermail/gcc-patches/2023-November/635176.html
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

3 participants