aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorTom de Vries <tdevries@suse.de>2020-09-09 09:51:43 +0200
committerTom de Vries <tdevries@suse.de>2020-09-09 14:33:19 +0200
commit505590b796df18ec3fcdcd6b8060f6f1410660b2 (patch)
tree6c238eb226a910cd607058b58d758dad5603d787
parentdb918db2c305adf8ba727a5981a684c19de1510e (diff)
downloadgcc-505590b796df18ec3fcdcd6b8060f6f1410660b2.zip
gcc-505590b796df18ec3fcdcd6b8060f6f1410660b2.tar.gz
gcc-505590b796df18ec3fcdcd6b8060f6f1410660b2.tar.bz2
[nvptx] Fix boolean type test in write_fn_proto
When running this libgomp testcase for nvptx accelerator: ... /* { dg-do run } */ __uint128_t v; int main () { #pragma omp target { __uint128_t exp = 2; __atomic_compare_exchange_n (&v, &exp, 7, false, __ATOMIC_RELEASE, __ATOMIC_ACQUIRE); } } ... we run into this assert in write_fn_proto: ... 913 gcc_assert (type == boolean_type_node); ... This happens when doing some special-handling code for __atomic_compare_exchange_1/2/4/8/16. The function decls have a parameter called weak of type bool, which is skipped when writing the decl because the corresponding libatomic functions do not have that parameter. The assert is there to verify that we skip the correct parameter. However, we assert because we have different type of bools: ... (gdb) call debug_generic_expr (type) _Bool (gdb) call debug_generic_expr (global_trees[TI_BOOLEAN_TYPE]) bool ... Fix this by checking for TREE_CODE (type) == BOOLEAN_TYPE instead. Tested libgomp on x86_64-linux with nvptx accelerator. Likewise, tested that the test-case above does not ICE anymore. gcc/ChangeLog: PR target/96991 * config/nvptx/nvptx.c (write_fn_proto): Fix boolean type check.
-rw-r--r--gcc/config/nvptx/nvptx.c2
1 files changed, 1 insertions, 1 deletions
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 39d0275..6f393df 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -910,7 +910,7 @@ write_fn_proto (std::stringstream &s, bool is_defn,
if (not_atomic_weak_arg)
argno = write_arg_type (s, -1, argno, type, prototyped);
else
- gcc_assert (type == boolean_type_node);
+ gcc_assert (TREE_CODE (type) == BOOLEAN_TYPE);
}
if (stdarg_p (fntype))