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