[gcc(refs/vendors/redhat/heads/gcc-8-branch)] [HSA] Avoid ICE when "HSA does not implement indirect calls"

[Jakub Jelinek <jakub@gcc.gnu.org>]

https://gcc.gnu.org/g:bb15082628a75bb3dfe2fbf5cdb33847ee22e324

commit bb15082628a75bb3dfe2fbf5cdb33847ee22e324
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Sat Jun 6 18:44:34 2020 +0200

    [HSA] Avoid ICE when "HSA does not implement indirect calls"
    
    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=<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
    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.
    
    (cherry picked from commit 973bce0fb50bbfd91f47238b82b99935525716ad)

Diff:
---
 gcc/hsa-gen.c                                      | 11 +++++-----
 .../c-c++-common/gomp/hsa-indirect-call-1.c        | 24 ++++++++++++++++++++++
 2 files changed, 30 insertions(+), 5 deletions(-)

diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 7974fffe360..5a4b38d717b 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -5251,11 +5251,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 (DECL_BUILT_IN_CLASS (function_decl)
-	  && DECL_FUNCTION_CODE (function_decl) == BUILT_IN_PREFETCH)
-	return;
 
       if (function_decl == NULL_TREE)
 	{
@@ -5264,6 +5259,12 @@ 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 (DECL_BUILT_IN_CLASS (function_decl)
+	  && DECL_FUNCTION_CODE (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 } */