procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

[HSA] Avoid ICE when "HSA does not implement indirect calls"

Browse source 
Made apparent by recent commit dc703151d4
"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=<optimized out>) 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
56b1c60e41 (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.
This commit is contained in:
Thomas Schwinge 2020-06-06 18:44:34 +02:00
parent a4df94a033
commit 973bce0fb5
2 changed files with 29 additions and 4 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

9
gcc/hsa-gen.c
View file

@ -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))

24
gcc/testsuite/c-c++-common/gomp/hsa-indirect-call-1.c Normal file
View file

@ -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 } */