[gomp4] [WIP] OpenACC bind, nohost clauses

Cesar Philippidis cesar_philippidis@mentor.com
Wed Dec 16 23:15:00 GMT 2015


On 12/14/2015 12:36 PM, Cesar Philippidis wrote:
> On 12/08/2015 11:55 AM, Thomas Schwinge wrote:
>> On Sat, 14 Nov 2015 09:36:36 +0100, I wrote:

>> C front end:
>>
>>     --- gcc/c/c-parser.c
>>     +++ gcc/c/c-parser.c
>>     @@ -11607,6 +11607,8 @@ c_parser_oacc_clause_async (c_parser *parser, tree list)
>>      static tree
>>      c_parser_oacc_clause_bind (c_parser *parser, tree list)
>>      {
>>     +  check_no_duplicate_clause (list, OMP_CLAUSE_BIND, "bind");
>>     +
>>        location_t loc = c_parser_peek_token (parser)->location;
>>      
>>        parser->lex_untranslated_string = true;
>>     @@ -11615,20 +11617,43 @@ c_parser_oacc_clause_bind (c_parser *parser, tree list)
>>            parser->lex_untranslated_string = false;
>>            return list;
>>          }
>>     -  if (c_parser_next_token_is (parser, CPP_NAME)
>>     -      || c_parser_next_token_is (parser, CPP_STRING))
>>     +  tree name = error_mark_node;
>>     +  c_token *token = c_parser_peek_token (parser);
>>     +  if (c_parser_next_token_is (parser, CPP_NAME))
>>          {
>>     -      tree t = c_parser_peek_token (parser)->value;
>>     +      tree decl = lookup_name (token->value);
>>     +      if (!decl)
>>     +       error_at (token->location, "%qE has not been declared",
>>     +                 token->value);
>>     +      else if (TREE_CODE (decl) != FUNCTION_DECL)
>>     +       error_at (token->location, "%qE does not refer to a function",
>>     +                 token->value);
>>
>> Quite possibly we'll want to add more error checking (matching signature
>> of X and Y, for example).
> 
> Good idea, but I wonder if that would be too strict. Should we allow
> integer promotion in the bind function arguments?

I decided to be strict here. In c++ there is a decls_match function
which determines if, say, a function prototype matches it's definition.
Turns out that function was a little too strict and generic for the bind
clause, because I think we want to allow the user to bind functions
declared in different namespaces. E.g.

 namespace foo {
   #pragma acc routine
   int bar ();
 }

 #pragma acc routine bind (foo::bar)
 ...

This should be acceptable. As a consequence, I created a
bind_decls_match function, for this purpose. I probably could have
taught the existing decls_match function how to copy with bind clauses,
but I suspect we may need to add more special cases for the bind clause
because the spec is so vague.

>> Again simplifying the c_head/clauses handling (snipped), the C++ front
>> end changes are very similar to the C front end changes:
>>
>>     --- gcc/cp/parser.c
>>     +++ gcc/cp/parser.c
>>     @@ -31539,42 +31538,76 @@ static tree
>>      cp_parser_oacc_clause_bind (cp_parser *parser, tree list)
>>      {
>>     [...]
>>     -  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
>>     -      || cp_lexer_next_token_is (parser->lexer, CPP_STRING))
>>     +  tree name = error_mark_node;
>>     +  cp_token *token = cp_lexer_peek_token (parser->lexer);
>>     +  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
>>
>> I'm not particularly confident in the following lookup/error checking
>> (which I copied a lot from C++ OpenACC routine parsing):
>>
>>          {
>>     -      tree t;
>>     -
>>     -      if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING)
>>     -       {
>>     -         t = cp_lexer_peek_token (parser->lexer)->u.value;
>>     -         cp_lexer_consume_token (parser->lexer);
>>     +      //TODO
>>     +      tree id = cp_parser_id_expression (parser, /*template_p=*/false,
>>     +                                        /*check_dependency_p=*/true,
>>     +                                        /*template_p=*/NULL,
>>     +                                        /*declarator_p=*/false,
>>     +                                        /*optional_p=*/false);
>>     +      tree decl = cp_parser_lookup_name_simple (parser, id, token->location);
>>     +      if (id != error_mark_node && decl == error_mark_node)
>>     +       cp_parser_name_lookup_error (parser, id, decl, NLE_NULL,
>>     +                                    token->location);
>>     +      if (/* TODO */ !decl || decl == error_mark_node)
>>     +       error_at (token->location, "%qE has not been declared",
>>     +                 token->u.value);
>>     +      else if (/* TODO */ is_overloaded_fn (decl)
>>     +              && (TREE_CODE (decl) != FUNCTION_DECL
>>     +                  || DECL_FUNCTION_TEMPLATE_P (decl)))
>>     +       error_at (token->location, "%qE names a set of overloads",
>>     +                 token->u.value);
>>     +      else if (/* TODO */ !DECL_NAMESPACE_SCOPE_P (decl))
>>     +       {
>>     +         /* Perhaps we should use the same rule as declarations in different
>>     +            namespaces?  */
>>     +         error_at (token->location,
>>     +                   "%qE does not refer to a namespace scope function",
>>     +                   token->u.value);
>>             }

This namespace check is happening too early, and I don't think it's
necessary either.

>>     +      else if (TREE_CODE (decl) != FUNCTION_DECL)
>>     +       error_at (token->location,
>>     +                 "%qE does not refer to a function",
>>     +                 token->u.value);
>>
>> ... also we'll want to add a lot more testsuite coverage for this.  (Also
>> for the OpenACC routine directive itself.)

I added some more test coverage, and I found and fixed a minor bug in
cp_finalize_oacc_routine.

>>            else
>>     -       t = cp_parser_id_expression (parser, /*template_p=*/false,
>>     -                                    /*check_dependency_p=*/true,
>>     -                                    /*template_p=*/NULL,
>>     -                                    /*declarator_p=*/false,
>>     -                                    /*optional_p=*/false);
>>     -      if (t == error_mark_node)
>>     -       return t;
>>     -
>>     -      tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
>>     -      OMP_CLAUSE_BIND_NAME (c) = t;
>>     -      OMP_CLAUSE_CHAIN (c) = list;
>>     -      list = c;
>>     +       {
>>     +         //TODO? TREE_USED (decl) = 1;
>>     +         tree name_id = DECL_NAME (decl);
>>     +         name = build_string (IDENTIFIER_LENGTH (name_id),
>>     +                              IDENTIFIER_POINTER (name_id));
>>
>> We probably need to apply C++ name mangling here?  How to do that?

Yes. But I deferred this to cp_finalize_oacc_routine because
cp_parser_oacc_clause_bind doesn't have access to the function decl.
Later on, the function decl gets mangled by calling decl_assembler_name.

>>     +       }
>>     +      //cp_lexer_consume_token (parser->lexer);
>>     +    }
>>     +  else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING))
>>     +    {
>>     +      name = token->u.value;
>>     +      cp_lexer_consume_token (parser->lexer);

I added a new check over here for empty strings, e.g. bind (""). Are
there any other obviously invalid strings?

>>          }
>>        else
>>     -    cp_parser_error (parser, "expected identifier or character string literal");
>>     +    cp_parser_error (parser,
>>     +                    "expected identifier or character string literal");
>>        parser->translate_strings_p = save_translate_strings_p;
>>        cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN);
>>     +  if (name != error_mark_node)
>>     +    {
>>     +      tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
>>     +      OMP_CLAUSE_BIND_NAME (c) = name;
>>     +      OMP_CLAUSE_CHAIN (c) = list;
>>     +      list = c;
>>     +    }
>>        return list;
>>      }
>>     --- gcc/omp-low.c
>>     +++ gcc/omp-low.c
>>     @@ -19853,14 +19857,76 @@ default_goacc_reduction (gcall *call)
>>      static unsigned int
>>      execute_oacc_device_lower ()
>>      {
>>     -  tree attrs = get_oacc_fn_attrib (current_function_decl);
>>     -  int dims[GOMP_DIM_MAX];
>>     -  
>>     -  if (!attrs)
>>     +  /* There are offloaded functions without an "omp declare target" attribute,
>>     +     so we'll not handle these here, but on the other hand, OpenACC bind and
>>     +     nohost clauses can only be generated in the front ends, and an "omp
>>     +     declare target" attribute will then also always have been set there, so
>>     +     this is not a problem in practice.  */
>>     +  tree attr = lookup_attribute ("omp declare target",
>>     +                               DECL_ATTRIBUTES (current_function_decl));
>>     +
>>     +#if defined(ACCEL_COMPILER)
>>     +  /* In an offload compiler, discard any offloaded function X that is tagged
>>     +     with an OpenACC bind(Y) clause: all references to X have been rewritten to
>>     +     refer to Y; X is unreachable, do not compile it.  */
>>     +  if (attr)
>>     +    {
>>     +      tree clauses = TREE_VALUE (attr);
>>     +      /* TODO: device_type handling.  */
>>     +      tree clause_bind = find_omp_clause (clauses, OMP_CLAUSE_BIND);
>>     +      if (clause_bind)
>>     +       {
>>     +         tree clause_bind_name = OMP_CLAUSE_BIND_NAME (clause_bind);
>>     +         const char *bind_name = TREE_STRING_POINTER(clause_bind_name);
>>     +         if (dump_file)
>>     +           fprintf (dump_file,
>>     +                    "Discarding function \"%s\" with \"bind(%s)\" clause.\n",
>>     +                    IDENTIFIER_POINTER (DECL_NAME (current_function_decl)),
>>     +                    bind_name);
>>     +         TREE_ASM_WRITTEN (current_function_decl) = 1;
>>     +         return TODO_discard_function;
>>     +       }
>>     +    }
>>     +#endif /* ACCEL_COMPILER */
>>     +#if !defined(ACCEL_COMPILER)
>>     +  /* In the host compiler, discard any offloaded function that is tagged with
>>     +     an OpenACC nohost clause.  */
>>     +  if (attr)
>>     +    {
>>     +      tree clauses = TREE_VALUE (attr);
>>     +      if (find_omp_clause (clauses, OMP_CLAUSE_NOHOST))
>>     +       {
>>     +         /* There are no construct/clause combinations that could make this
>>     +            happen, but play it safe, and verify that we never discard a
>>     +            function that is stored in offload_funcs, used for target/offload
>>     +            function mapping.  */
>>     +         if (flag_checking)
>>     +           {
>>     +             bool found = false;
>>     +             for (unsigned i = 0;
>>     +                  !found && i < vec_safe_length (offload_funcs);
>>     +                  i++)
>>     +               if ((*offload_funcs)[i] == current_function_decl)
>>     +                 found = true;
>>     +             gcc_assert (!found);
>>     +           }
>>     +
>>     +         if (dump_file)
>>     +           fprintf (dump_file,
>>     +                    "Discarding function \"%s\" with \"nohost\" clause.\n",
>>     +                    IDENTIFIER_POINTER (DECL_NAME (current_function_decl)));
>>     +         TREE_ASM_WRITTEN (current_function_decl) = 1;
>>     +         return TODO_discard_function;
> 
> I don't think this is a good idea. If you have a nohost function,
> wounldn't that prevent the code from linking?
> 
> Perhaps nohost should kind of implement a reverse bind on the host. E.g.
> discard the function defintion and replace it with an asm alias to some
> libgomp function like goacc_nohost_fallback. That way, the program will
> still link and the runtime will provide the end user with a sensible
> error when things go wrong.

Are you planning on working on this?

Note that my test coverage is limited to compile-only tests for now. I'd
like to test nohost and bind at runtime, but at least the nohosts tests
depend on this change. Do you want to work on the runtime tests?

I'll apply this patch to gomp4 in a couple of days. In the meantime,
glance at the error messages and see if they are sufficiently clear. Or
maybe I'm missing some test coverage.

Cesar

-------------- next part --------------
A non-text attachment was scrubbed...
Name: cxx-bind-20151216.diff
Type: text/x-patch
Size: 13193 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20151216/479e6cc0/attachment.bin>


More information about the Gcc-patches mailing list