From 973bce0fb50bbfd91f47238b82b99935525716ad Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Sat, 6 Jun 2020 18:44:34 +0200 Subject: [PATCH] [HSA] Avoid ICE when "HSA does not implement indirect calls" MIME-Version: 1.0 Content-Type: text/plain; charset=utf8 Content-Transfer-Encoding: 8bit Made apparent by recent commit dc703151d4f4560e647649506d5b4ceb0ee11e90 "openmp: Implement discovery of implicit declare target to clauses": +FAIL: libgomp.c/target-39.c (internal compiler error) +FAIL: libgomp.c/target-39.c (test for excess errors) +UNRESOLVED: libgomp.c/target-39.c compilation failed to produce executable This is in a '--enable-offload-targets=[...],hsa' build, with '-foffload=hsa' enabled (by default). during GIMPLE pass: hsagen source-gcc/libgomp/testsuite/libgomp.c/target-39.c: In function ‘main._omp_fn.0.hsa.0’: source-gcc/libgomp/testsuite/libgomp.c/target-39.c:23:11: internal compiler error: Segmentation fault 23 | #pragma omp target map(from:err) | ^~~ [...] GDB: Program received signal SIGSEGV, Segmentation fault. fndecl_built_in_p (node=0x0, name=BUILT_IN_PREFETCH) at [...]/source-gcc/gcc/tree.h:6267 6267 return (fndecl_built_in_p (node, BUILT_IN_NORMAL) (gdb) bt #0 fndecl_built_in_p (node=0x0, name=BUILT_IN_PREFETCH) at [...]/source-gcc/gcc/tree.h:6267 #1 0x0000000000b19739 in gen_hsa_insns_for_call (stmt=stmt@entry=0x7ffff693b200, hbb=hbb@entry=0x2b152c0) at [...]/source-gcc/gcc/hsa-gen.c:5304 #2 0x0000000000b1aca7 in gen_hsa_insns_for_gimple_stmt (stmt=0x7ffff693b200, hbb=hbb@entry=0x2b152c0) at [...]/source-gcc/gcc/hsa-gen.c:5770 #3 0x0000000000b1bd21 in gen_body_from_gimple () at [...]/source-gcc/gcc/hsa-gen.c:5999 #4 0x0000000000b1dbd2 in generate_hsa (kernel=) at [...]/source-gcc/gcc/hsa-gen.c:6596 #5 0x0000000000b1de66 in (anonymous namespace)::pass_gen_hsail::execute (this=0x2a2aac0) at [...]/source-gcc/gcc/hsa-gen.c:6680 #6 0x0000000000d06f90 in execute_one_pass (pass=pass@entry=0x2a2aac0) at [...]/source-gcc/gcc/passes.c:2502 [...] (gdb) up #1 0x0000000000b19739 in gen_hsa_insns_for_call (stmt=stmt@entry=0x7ffff693b200, hbb=hbb@entry=0x2b152c0) at /home/thomas/tmp/source/gcc/build/track-slim-omp/source-gcc/gcc/hsa-gen.c:5304 5304 if (fndecl_built_in_p (function_decl, BUILT_IN_PREFETCH)) (gdb) print function_decl $1 = (tree) 0x0 (gdb) list 5299 if (!gimple_call_builtin_p (stmt, BUILT_IN_NORMAL)) 5300 { 5301 tree function_decl = gimple_call_fndecl (stmt); 5302 /* Prefetch pass can create type-mismatching prefetch builtin calls which 5303 fail the gimple_call_builtin_p test above. Handle them here. */ 5304 if (fndecl_built_in_p (function_decl, BUILT_IN_PREFETCH)) 5305 return; 5306 5307 if (function_decl == NULL_TREE) 5308 { The problem is present already since 2016-11-23 commit 56b1c60e412fcf1245b4780871553cbdebb956a3 (r242761) "Merge from HSA branch to trunk", and the fix obvious enough. gcc/ * hsa-gen.c (gen_hsa_insns_for_call): Move 'function_decl == NULL_TREE' check earlier. gcc/testsuite/ * c-c++-common/gomp/hsa-indirect-call-1.c: New file. --- gcc/hsa-gen.c | 9 +++---- .../c-c++-common/gomp/hsa-indirect-call-1.c | 24 +++++++++++++++++++ 2 files changed, 29 insertions(+), 4 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/hsa-indirect-call-1.c diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index 43baa2e82c8..2af999048b2 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -5299,10 +5299,6 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb) if (!gimple_call_builtin_p (stmt, BUILT_IN_NORMAL)) { tree function_decl = gimple_call_fndecl (stmt); - /* Prefetch pass can create type-mismatching prefetch builtin calls which - fail the gimple_call_builtin_p test above. Handle them here. */ - if (fndecl_built_in_p (function_decl, BUILT_IN_PREFETCH)) - return; if (function_decl == NULL_TREE) { @@ -5311,6 +5307,11 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb) return; } + /* Prefetch pass can create type-mismatching prefetch builtin calls which + fail the gimple_call_builtin_p test above. Handle them here. */ + if (fndecl_built_in_p (function_decl, BUILT_IN_PREFETCH)) + return; + if (hsa_callable_function_p (function_decl)) gen_hsa_insns_for_direct_call (stmt, hbb); else if (!gen_hsa_insns_for_known_library_call (stmt, hbb)) diff --git a/gcc/testsuite/c-c++-common/gomp/hsa-indirect-call-1.c b/gcc/testsuite/c-c++-common/gomp/hsa-indirect-call-1.c new file mode 100644 index 00000000000..67ee6af309a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/hsa-indirect-call-1.c @@ -0,0 +1,24 @@ +/* Instead of ICE, we'd like "HSA does not implement indirect calls". */ + +/* Reduced from 'libgomp.c/target-39.c'. */ + +/* { dg-require-effective-target offload_hsa } */ +/* { dg-additional-options "-Whsa" } to override '{gcc,g++}.dg/gomp/gomp.exp'. */ + +typedef void (*fnp) (void); +void f1 (void) { } +fnp f2 (void) { return f1; } +#pragma omp declare target to (f1, f2) + +int +main () +{ + #pragma omp target + { + fnp fnp = f2 (); + fnp (); /* { dg-message "note: support for HSA does not implement indirect calls" } */ + } + return 0; +} + +/* { dg-warning "could not emit HSAIL for the function" "" { target *-*-* } 0 } */ -- 2.30.2