From patchwork Sat May 4 21:21:45 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 1931426 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=VB7nLLX+; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4VX10C2zhzz1ybC for ; Sun, 5 May 2024 07:24:43 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9ACB4386F434 for ; Sat, 4 May 2024 21:24:41 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-il1-x130.google.com (mail-il1-x130.google.com [IPv6:2607:f8b0:4864:20::130]) by sourceware.org (Postfix) with ESMTPS id 2E1FB384476F for ; Sat, 4 May 2024 21:22:11 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 2E1FB384476F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 2E1FB384476F Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2607:f8b0:4864:20::130 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1714857737; cv=none; b=RWQk7CFsXE0D17KOXcgJhvADrI5oUGA4u/e8ATLWOXJFqzquSu07AD6sPVv8/oJeVsxrwQGIqApHAGzJ3lyYIoOl1HPQ0L6DyC7148odGtnr/u8WDzTgjCRHA+S+ma44m504NNOYZycS8/3fJZ84Pm2JmuphnXTOx/+g9wjnvHU= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1714857737; c=relaxed/simple; bh=q0NTDWtU2jdqJxn3QpO2Q5ZUdwccFoz+6mx7co7gN9A=; h=DKIM-Signature:From:To:Subject:Date:Message-Id:MIME-Version; b=MocfgGonTRH/PV9rdd9E8lGutWD/vZCPYXAPIl5UAj5pIo9GqXdM8Ji74TLn3cqDin1KXQUxSN63YwIo7HceIIuu6UIrmJCCwfjoVjmfcvIMEjF4QGm1jKhXlipFvgRTgoZMoHydPstJRP6few/oxbafSX3oRBZUUfiAgn/IZNo= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-il1-x130.google.com with SMTP id e9e14a558f8ab-36c82ca80adso4079965ab.3 for ; Sat, 04 May 2024 14:22:11 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1714857729; x=1715462529; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=sDkDHgtBmQ8sixNL4cbsof8xJVJJPfFcuzgogLJij10=; b=VB7nLLX+zsBIb0OjCZyog5nQoJWNf3Bmby02vi3AqjC6wFly8tdqTKmPDzclB89JhJ OSf+0X7qu9JPV81WGMXl4GENtlGtVxV0L6+8QLXMRHod/jmJFcyoiznbtdU6+UEtMN/2 nkau6UD708LI8jexEVQ3drd91HT0bxASbq4oSJA77Kb4aO8UMBTzpKmkdFoTQg5UtpNq vgJIiBeniLarf+Qky9H7/gFColmfchgqTQnIjxERY/n6mwiq0V9iCyiNpn5S3zFSIHvq mva598MGQ+OHuRHvYUSIzDZjp5721CmXFoOy7Jio/ezuFgyztxM82t3DZ9F/6thmPQDS m4Fg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1714857729; x=1715462529; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to; bh=sDkDHgtBmQ8sixNL4cbsof8xJVJJPfFcuzgogLJij10=; b=Da67Vte4krG7iHLI2brsHt8gsfz4VJ8V5Ik7eiRulgLPwquxXGr+q4ALHVkHmca2Vu ScU4h49KJel5+DCVw3aQ6Ha55Z2e5BGBwwfe58ihqxDo6qZe6ujNkxmbSCXcTh4bVRQJ 5fb/iQ8nXVPXEXk8oe9FpZCSSmAG3HzptNrEtdfIWaV0dwHod3tEU2DEFBl/O+8aPdG1 9ywfMepqivvfieCrjC0JSbQcm5uABzP99U+BrFyfqzXM5ZyE91wRJwA013ASvvLlNdnE fIaolDstttz5q65R3T4n/fBhfS3ioR98r1pxWVtHGlTXDIUE/rgds/0WQ+Jy7nzJo9s9 Y1UQ== X-Gm-Message-State: AOJu0YxbOYCu+QoFOsZYaEHW8ljvfIRZybpSvgb3jeSDfasjfiHBgKQ0 h8xQgbOL0MZm9onGOaQlTm+dJBGP+4/MaPrt6S2l/EveFU2VYSTaWlWVRvigLXyF4KgWC+E7yVi 9 X-Google-Smtp-Source: AGHT+IE0zYOoeRKw4k13ALPu9qJcF4vQ4EDIFsms0Jm2quPMCEV3x6vvwhGTYGIWb9aOQ86j3+TwUQ== X-Received: by 2002:a05:6e02:16cd:b0:36c:a4e:9ea8 with SMTP id 13-20020a056e0216cd00b0036c0a4e9ea8mr8487287ilx.17.1714857729082; Sat, 04 May 2024 14:22:09 -0700 (PDT) Received: from pondscum.hsd1.co.comcast.net ([2601:281:d901:5620:3e29:4728:ec99:5098]) by smtp.gmail.com with ESMTPSA id ez3-20020a056638614300b004877be21febsm1559468jab.62.2024.05.04.14.22.08 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Sat, 04 May 2024 14:22:08 -0700 (PDT) From: Sandra Loosemore To: gcc-patches@gcc.gnu.org Cc: jakub@redhat.com, tburnus@baylibre.com Subject: [PATCH 05/12] OpenMP: C++ front-end support for metadirectives Date: Sat, 4 May 2024 15:21:45 -0600 Message-Id: <20240504212153.3561429-6-sloosemore@baylibre.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20240504212153.3561429-1-sloosemore@baylibre.com> References: <20240504212153.3561429-1-sloosemore@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-11.1 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org This patch adds C++ support for metadirectives. It uses the c-family support committed with the corresponding C front end patch to do early parse-time metadirective resolution when possible. Additional C/C++ common testcases are provided in a subsequent patch in the series. gcc/cp/ChangeLog * parser.cc (cp_parser_skip_to_end_of_block_or_statement): Add metadirective_p parameter, use it to control brace/parentheses behavior for metadirectives. (mangle_metadirective_region_label): New. (cp_parser_label_for_labeled_statement): Use it. (cp_parser_jump_statement): Likewise. (cp_parser_omp_context_selector): Add metadirective_p parameter, use it to control error behavior for non-constant exprs properties. (cp_parser_omp_context_selector_specification): Add metadirective_p parameter, use it for cp_parser_omp_context_selector call. (cp_finish_omp_declare_variant): Adjust call to cp_parser_omp_context_selector_specification. (analyze_metadirective_body): New. (cp_parser_omp_metadirective): New. (cp_parser_pragma): Handle PRAGMA_OMP_METADIRECTIVE. * parser.h (struct cp_parser): Add fields for metadirective parsing state. * pt.cc (tsubst_omp_context_selector): New. (tsubst_stmt): Handle OMP_METADIRECTIVE. gcc/testsuite/ChangeLog * g++.dg/gomp/attrs-metadirective-1.C: New. * g++.dg/gomp/attrs-metadirective-2.C: New. * g++.dg/gomp/attrs-metadirective-3.C: New. * g++.dg/gomp/attrs-metadirective-4.C: New. * g++.dg/gomp/attrs-metadirective-5.C: New. * g++.dg/gomp/attrs-metadirective-6.C: New. * g++.dg/gomp/attrs-metadirective-7.C: New. * g++.dg/gomp/attrs-metadirective-8.C: New. libgomp/ChangeLog * testsuite/libgomp.c++/metadirective-template-1.C: New. * testsuite/libgomp.c++/metadirective-template-2.C: New. * testsuite/libgomp.c++/metadirective-template-3.C: New. Co-Authored-By: Kwok Cheung Yeung Co-Authored-By: Sandra Loosemore --- gcc/cp/parser.cc | 524 +++++++++++++++++- gcc/cp/parser.h | 7 + gcc/cp/pt.cc | 119 ++++ .../g++.dg/gomp/attrs-metadirective-1.C | 40 ++ .../g++.dg/gomp/attrs-metadirective-2.C | 74 +++ .../g++.dg/gomp/attrs-metadirective-3.C | 31 ++ .../g++.dg/gomp/attrs-metadirective-4.C | 41 ++ .../g++.dg/gomp/attrs-metadirective-5.C | 24 + .../g++.dg/gomp/attrs-metadirective-6.C | 31 ++ .../g++.dg/gomp/attrs-metadirective-7.C | 31 ++ .../g++.dg/gomp/attrs-metadirective-8.C | 16 + .../libgomp.c++/metadirective-template-1.C | 37 ++ .../libgomp.c++/metadirective-template-2.C | 41 ++ .../libgomp.c++/metadirective-template-3.C | 41 ++ 14 files changed, 1044 insertions(+), 13 deletions(-) create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C create mode 100644 libgomp/testsuite/libgomp.c++/metadirective-template-1.C create mode 100644 libgomp/testsuite/libgomp.c++/metadirective-template-2.C create mode 100644 libgomp/testsuite/libgomp.c++/metadirective-template-3.C diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 8b819b2ebfd..4bb9b086095 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -3001,7 +3001,7 @@ static void cp_parser_skip_to_end_of_statement static void cp_parser_consume_semicolon_at_end_of_statement (cp_parser *); static void cp_parser_skip_to_end_of_block_or_statement - (cp_parser *); + (cp_parser *, bool = false); static bool cp_parser_skip_to_closing_brace (cp_parser *); static bool cp_parser_skip_entire_template_parameter_list @@ -4177,9 +4177,11 @@ cp_parser_consume_semicolon_at_end_of_statement (cp_parser *parser) have consumed a non-nested `;'. */ static void -cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser) +cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser, + bool metadirective_p) { int nesting_depth = 0; + int bracket_depth = 0; /* Unwind generic function template scope if necessary. */ if (parser->fully_implicit_function_template_p) @@ -4201,7 +4203,7 @@ cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser) case CPP_SEMICOLON: /* Stop if this is an unnested ';'. */ - if (!nesting_depth) + if (!nesting_depth && (!metadirective_p || bracket_depth <= 0)) nesting_depth = -1; break; @@ -4220,6 +4222,19 @@ cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser) nesting_depth++; break; + case CPP_OPEN_PAREN: + /* Track parentheses in case the statement is a standalone 'for' + statement - we want to skip over the semicolons separating the + operands. */ + if (metadirective_p && nesting_depth == 0) + bracket_depth++; + break; + + case CPP_CLOSE_PAREN: + if (metadirective_p && nesting_depth == 0) + bracket_depth--; + break; + case CPP_KEYWORD: if (!cp_token_is_module_directive (token)) break; @@ -12999,6 +13014,18 @@ attr_chainon (tree attrs, tree attr) return chainon (attrs, attr); } + +/* Helper function for cp_parser_label: mangle a metadirective region + label NAME. */ +static tree +mangle_metadirective_region_label (cp_parser *parser, tree name) +{ + const char *old_name = IDENTIFIER_POINTER (name); + char *new_name = (char *) alloca (strlen (old_name) + 32); + sprintf (new_name, "%s_MDR%u", old_name, parser->metadirective_region_num); + return get_identifier (new_name); +} + /* Parse the label for a labeled-statement, i.e. label: @@ -13101,7 +13128,12 @@ cp_parser_label_for_labeled_statement (cp_parser* parser, tree attributes) default: /* Anything else must be an ordinary label. */ - label = finish_label_stmt (cp_parser_identifier (parser)); + cp_expr identifier = cp_parser_identifier (parser); + if (identifier != error_mark_node + && parser->in_metadirective_body + && parser->metadirective_body_labels->contains (*identifier)) + *identifier = mangle_metadirective_region_label (parser, *identifier); + label = finish_label_stmt (identifier); if (label && TREE_CODE (label) == LABEL_DECL) FALLTHROUGH_LABEL_P (label) = fallthrough_p; break; @@ -14906,7 +14938,15 @@ cp_parser_jump_statement (cp_parser* parser) finish_goto_stmt (cp_parser_expression (parser)); } else - finish_goto_stmt (cp_parser_identifier (parser)); + { + cp_expr identifier = cp_parser_identifier (parser); + if (identifier != error_mark_node + && parser->in_metadirective_body + && parser->metadirective_body_labels->contains (*identifier)) + *identifier = mangle_metadirective_region_label (parser, + *identifier); + finish_goto_stmt (identifier); + } /* Look for the final `;'. */ cp_parser_require (parser, CPP_SEMICOLON, RT_SEMICOLON); break; @@ -47899,7 +47939,7 @@ cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok, static tree cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set, - bool has_parms_p) + bool has_parms_p, bool metadirective_p) { tree ret = NULL_TREE; do @@ -48060,17 +48100,25 @@ cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set, break; case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR: case OMP_TRAIT_PROPERTY_BOOL_EXPR: - /* FIXME: this is bogus, the expression need - not be constant. */ - t = cp_parser_constant_expression (parser); + /* FIXME: I believe it is an unimplemented feature rather + than a user error to have non-constant expressions + inside "declare variant". */ + t = metadirective_p + ? cp_parser_expression (parser) + : cp_parser_constant_expression (parser); if (t != error_mark_node) { t = fold_non_dependent_expr (t); - if (!value_dependent_expression_p (t) + if (!metadirective_p + && !value_dependent_expression_p (t) && (!INTEGRAL_TYPE_P (TREE_TYPE (t)) || !tree_fits_shwi_p (t))) error_at (token->location, "property must be " "constant integer expression"); + if (metadirective_p + && !INTEGRAL_TYPE_P (TREE_TYPE (t))) + error_at (token->location, + "property must be integer expression"); else properties = make_trait_property (NULL_TREE, t, properties); @@ -48149,7 +48197,8 @@ cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set, static tree cp_parser_omp_context_selector_specification (cp_parser *parser, - bool has_parms_p) + bool has_parms_p, + bool metadirective_p) { tree ret = NULL_TREE; do @@ -48176,7 +48225,8 @@ cp_parser_omp_context_selector_specification (cp_parser *parser, return error_mark_node; tree selectors - = cp_parser_omp_context_selector (parser, set, has_parms_p); + = cp_parser_omp_context_selector (parser, set, has_parms_p, + metadirective_p); if (selectors == error_mark_node) { cp_parser_skip_to_closing_brace (parser); @@ -48486,7 +48536,8 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok, if (!parens.require_open (parser)) goto fail; - tree ctx = cp_parser_omp_context_selector_specification (parser, true); + tree ctx = cp_parser_omp_context_selector_specification (parser, true, + false); if (ctx == error_mark_node) goto fail; ctx = omp_check_context_selector (match_loc, ctx, false); @@ -49119,6 +49170,449 @@ cp_parser_omp_end (cp_parser *parser, cp_token *pragma_tok) } } + +/* Helper function for cp_parser_omp_metadirective. */ + +static void +analyze_metadirective_body (cp_parser *parser, + vec &tokens, + vec &labels) +{ + int nesting_depth = 0; + int bracket_depth = 0; + bool in_case = false; + bool in_label_decl = false; + cp_token *pragma_tok = NULL; + + while (1) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + bool stop = false; + + if (cp_lexer_next_token_is_keyword (parser->lexer, RID_CASE)) + in_case = true; + else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_LABEL)) + in_label_decl = true; + + switch (token->type) + { + case CPP_EOF: + break; + case CPP_NAME: + if ((!in_case + && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + || in_label_decl) + labels.safe_push (token->u.value); + goto add; + case CPP_OPEN_BRACE: + ++nesting_depth; + goto add; + case CPP_CLOSE_BRACE: + if (--nesting_depth == 0 && bracket_depth == 0) + stop = true; + goto add; + case CPP_OPEN_PAREN: + ++bracket_depth; + goto add; + case CPP_CLOSE_PAREN: + --bracket_depth; + goto add; + case CPP_COLON: + in_case = false; + goto add; + case CPP_SEMICOLON: + if (nesting_depth == 0 && bracket_depth == 0) + stop = true; + /* Local label declarations are terminated by a semicolon. */ + in_label_decl = false; + goto add; + case CPP_PRAGMA: + parser->lexer->in_pragma = true; + pragma_tok = token; + goto add; + case CPP_PRAGMA_EOL: + /* C++ attribute syntax for OMP directives lexes as a pragma, + but we must reset the associated lexer state when we reach + the end in order to get the tokens for the statement that + come after it. */ + tokens.safe_push (*token); + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + pragma_tok = NULL; + continue; + default: + add: + tokens.safe_push (*token); + cp_lexer_consume_token (parser->lexer); + if (stop) + break; + continue; + } + break; + } +} + +/* OpenMP 5.0: + + # pragma omp metadirective [clause[, clause]] +*/ + +static void +cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, + bool *if_p) +{ + static unsigned int metadirective_region_count = 0; + + auto_vec directive_tokens; + auto_vec body_tokens; + auto_vec body_labels; + auto_vec directives; + auto_vec ctxs; + bool default_seen = false; + int directive_token_idx = 0; + location_t pragma_loc = pragma_tok->location; + tree standalone_body = NULL_TREE; + vec candidates; + bool requires_body = false; + + tree ret = make_node (OMP_METADIRECTIVE); + SET_EXPR_LOCATION (ret, pragma_loc); + TREE_TYPE (ret) = void_type_node; + OMP_METADIRECTIVE_VARIANTS (ret) = NULL_TREE; + + while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) + { + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + if (cp_lexer_next_token_is_not (parser->lexer, CPP_NAME) + && cp_lexer_next_token_is_not (parser->lexer, CPP_KEYWORD)) + { + cp_parser_error (parser, "expected %, " + "%, or % clause"); + goto fail; + } + + location_t match_loc = cp_lexer_peek_token (parser->lexer)->location; + const char *p + = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); + cp_lexer_consume_token (parser->lexer); + bool default_p + = strcmp (p, "default") == 0 || strcmp (p, "otherwise") == 0; + if (default_p) + { + if (default_seen) + { + error_at (match_loc, "too many % or % " + "clauses in %"); + cp_parser_skip_to_end_of_block_or_statement (parser, true); + goto fail; + } + else + default_seen = true; + } + if (!strcmp (p, "when") == 0 && !default_p) + { + error_at (match_loc, "%qs is not valid for %qs", + p, "metadirective"); + cp_parser_skip_to_end_of_block_or_statement (parser, true); + goto fail; + } + + matching_parens parens; + tree ctx = NULL_TREE; + bool skip = false; + + if (!parens.require_open (parser)) + goto fail; + + if (!default_p) + { + ctx = cp_parser_omp_context_selector_specification (parser, false, + true); + if (ctx == error_mark_node) + goto fail; + ctx = omp_check_context_selector (match_loc, ctx, true); + if (ctx == error_mark_node) + goto fail; + + /* Remove the selector from further consideration if it can be + evaluated as a non-match at this point. */ + /* FIXME: we could still do this if the context selector + doesn't have any dependent subexpressions. */ + skip = (!processing_template_decl + && omp_context_selector_matches (ctx, true, true) == 0); + + if (cp_lexer_next_token_is_not (parser->lexer, CPP_COLON)) + { + cp_parser_require (parser, CPP_COLON, RT_COLON); + goto fail; + } + cp_lexer_consume_token (parser->lexer); + } + + /* Read in the directive type and create a dummy pragma token for + it. */ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + + const char *directive[3] = {}; + int i; + for (i = 0; i < 3; i++) + { + tree id; + if (cp_lexer_peek_nth_token (parser->lexer, i + 1)->type + == CPP_CLOSE_PAREN) + { + if (i == 0) + directive[i++] = "nothing"; + break; + } + else if (cp_lexer_peek_nth_token (parser->lexer, i + 1)->type + == CPP_NAME) + id = cp_lexer_peek_nth_token (parser->lexer, i + 1)->u.value; + else if (cp_lexer_peek_nth_token (parser->lexer, i + 1)->keyword + != RID_MAX) + { + enum rid rid + = cp_lexer_peek_nth_token (parser->lexer, i + 1)->keyword; + id = ridpointers[rid]; + } + else + break; + + directive[i] = IDENTIFIER_POINTER (id); + } + if (i == 0) + { + error_at (loc, "expected directive name"); + cp_parser_skip_to_end_of_block_or_statement (parser, true); + goto fail; + } + + const struct c_omp_directive *omp_directive + = c_omp_categorize_directive (directive[0], + directive[1], + directive[2]); + + if (omp_directive == NULL) + { + for (int j = 0; j < i; j++) + cp_lexer_consume_token (parser->lexer); + cp_parser_error (parser, "unknown directive name"); + goto fail; + } + else + { + int token_count = 0; + if (omp_directive->first) token_count++; + if (omp_directive->second) token_count++; + if (omp_directive->third) token_count++; + for (int j = 0; j < token_count; j++) + cp_lexer_consume_token (parser->lexer); + } + if (p == NULL) + { + cp_parser_error (parser, "expected directive name"); + goto fail; + } + if (omp_directive->id == PRAGMA_OMP_METADIRECTIVE) + { + cp_parser_error (parser, + "metadirectives cannot be used as variants of a " + "%"); + goto fail; + } + if (omp_directive->kind == C_OMP_DIR_DECLARATIVE) + { + sorry_at (loc, "declarative directive variants of a " + "% are not supported"); + goto fail; + } + if (omp_directive->kind == C_OMP_DIR_CONSTRUCT) + requires_body = true; + + if (!skip) + { + cp_token pragma_token; + pragma_token.type = CPP_PRAGMA; + pragma_token.location = loc; + pragma_token.u.value = build_int_cst (NULL, omp_directive->id); + + directives.safe_push (omp_directive); + directive_tokens.safe_push (pragma_token); + ctxs.safe_push (ctx); + } + + /* Read in tokens for the directive clauses. */ + int nesting_depth = 0; + while (1) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + switch (token->type) + { + case CPP_EOF: + case CPP_PRAGMA_EOL: + break; + case CPP_OPEN_PAREN: + ++nesting_depth; + goto add; + case CPP_CLOSE_PAREN: + if (nesting_depth-- == 0) + break; + goto add; + default: + add: + if (!skip) + directive_tokens.safe_push (*token); + cp_lexer_consume_token (parser->lexer); + continue; + } + break; + } + + cp_lexer_consume_token (parser->lexer); + + if (!skip) + { + cp_token eol_token = {}; + eol_token.type = CPP_PRAGMA_EOL; + eol_token.keyword = RID_MAX; + directive_tokens.safe_push (eol_token); + } + } + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + + if (!default_seen) + { + /* Add a default clause that evaluates to 'omp nothing'. */ + const struct c_omp_directive *omp_directive + = c_omp_categorize_directive ("nothing", NULL, NULL); + + cp_token pragma_token = {}; + pragma_token.type = CPP_PRAGMA; + pragma_token.keyword = RID_MAX; + pragma_token.location = UNKNOWN_LOCATION; + pragma_token.u.value = build_int_cst (NULL, PRAGMA_OMP_NOTHING); + + directives.safe_push (omp_directive); + directive_tokens.safe_push (pragma_token); + ctxs.safe_push (NULL_TREE); + + cp_token eol_token = {}; + eol_token.type = CPP_PRAGMA_EOL; + eol_token.keyword = RID_MAX; + directive_tokens.safe_push (eol_token); + } + + if (requires_body) + analyze_metadirective_body (parser, body_tokens, body_labels); + + /* Process each candidate directive. */ + unsigned i; + tree ctx; + cp_lexer *lexer; + + lexer = cp_lexer_alloc (); + lexer->debugging_p = parser->lexer->debugging_p; + vec_safe_reserve (lexer->buffer, + directive_tokens.length () + body_tokens.length () + 2); + + FOR_EACH_VEC_ELT (ctxs, i, ctx) + { + lexer->buffer->truncate (0); + + /* Add the directive tokens. */ + do + lexer->buffer->quick_push (directive_tokens [directive_token_idx++]); + while (lexer->buffer->last ().type != CPP_PRAGMA_EOL); + + /* Add the body tokens. */ + gcc_assert (requires_body || body_tokens.is_empty ()); + for (unsigned j = 0; j < body_tokens.length (); j++) + lexer->buffer->quick_push (body_tokens[j]); + + /* Make sure nothing tries to read past the end of the tokens. */ + cp_token eof_token = {}; + eof_token.type = CPP_EOF; + eof_token.keyword = RID_MAX; + lexer->buffer->quick_push (eof_token); + lexer->buffer->quick_push (eof_token); + + lexer->next_token = lexer->buffer->address(); + lexer->last_token = lexer->next_token + lexer->buffer->length () - 1; + + cp_lexer *old_lexer = parser->lexer; + bool old_in_metadirective_body = parser->in_metadirective_body; + vec *old_metadirective_body_labels + = parser->metadirective_body_labels; + unsigned int old_metadirective_region_num + = parser->metadirective_region_num; + parser->lexer = lexer; + cp_lexer_set_source_position_from_token (lexer->next_token); + parser->in_metadirective_body = true; + parser->metadirective_body_labels = &body_labels; + parser->metadirective_region_num = ++metadirective_region_count; + + int prev_errorcount = errorcount; + tree directive = push_stmt_list (); + tree directive_stmt = begin_compound_stmt (0); + + cp_parser_pragma (parser, pragma_compound, if_p); + finish_compound_stmt (directive_stmt); + directive = pop_stmt_list (directive); + + bool standalone_p + = directives[i]->kind == C_OMP_DIR_STANDALONE + || directives[i]->kind == C_OMP_DIR_UTILITY; + if (standalone_p && requires_body) + { + /* Parsing standalone directives will not consume the body + tokens, so do that here. */ + if (standalone_body == NULL_TREE) + { + standalone_body = push_stmt_list (); + cp_parser_statement (parser, NULL_TREE, false, if_p); + standalone_body = pop_stmt_list (standalone_body); + } + else + cp_parser_skip_to_end_of_block_or_statement (parser, true); + } + + tree body = standalone_p ? standalone_body : NULL_TREE; + tree variant = make_omp_metadirective_variant (ctx, directive, body); + OMP_METADIRECTIVE_VARIANTS (ret) + = chainon (OMP_METADIRECTIVE_VARIANTS (ret), variant); + + /* Check that all valid tokens have been consumed if no parse errors + encountered. */ + gcc_assert (errorcount != prev_errorcount + || cp_lexer_next_token_is (parser->lexer, CPP_EOF)); + + parser->lexer = old_lexer; + cp_lexer_set_source_position_from_token (old_lexer->next_token); + parser->in_metadirective_body = old_in_metadirective_body; + parser->metadirective_body_labels = old_metadirective_body_labels; + parser->metadirective_region_num = old_metadirective_region_num; + } + + /* Try to resolve the metadirective early. */ + if (!processing_template_decl) + { + candidates = omp_early_resolve_metadirective (ret); + if (!candidates.is_empty ()) + ret = c_omp_expand_metadirective (candidates); + } + + add_stmt (ret); + return; + +fail: + /* Skip the metadirective pragma. */ + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + + /* Skip the metadirective body. */ + cp_parser_skip_to_end_of_block_or_statement (parser, true); +} + + /* Helper function of cp_parser_omp_declare_reduction. Parse the combiner expression and optional initializer clause of #pragma omp declare reduction. We store the expression(s) as @@ -51051,6 +51545,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) cp_parser_omp_nothing (parser, pragma_tok); return false; + case PRAGMA_OMP_METADIRECTIVE: + cp_parser_omp_metadirective (parser, pragma_tok, if_p); + return true; + case PRAGMA_OMP_ERROR: return cp_parser_omp_error (parser, pragma_tok, context); diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h index 373e78f3ea4..1260898774c 100644 --- a/gcc/cp/parser.h +++ b/gcc/cp/parser.h @@ -446,6 +446,13 @@ struct GTY(()) cp_parser { /* Pointer to state for parsing omp_loops. Managed by cp_parser_omp_for_loop in parser.cc and not used outside that file. */ struct omp_for_parse_data * GTY((skip)) omp_for_parse_state; + + /* Set if we are processing a statement body associated with a + metadirective variant. */ + bool in_metadirective_body; + + vec * GTY((skip)) metadirective_body_labels; + unsigned metadirective_region_num; }; /* In parser.cc */ diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 3b2106dd3f6..409c4df68bc 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17859,6 +17859,79 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, return new_clauses; } +/* Like tsubst_copy, but specifically for OpenMP context selectors. */ +static tree +tsubst_omp_context_selector (tree ctx, tree args, tsubst_flags_t complain, + tree in_decl) +{ + tree new_ctx = NULL_TREE; + for (tree set = ctx; set; set = TREE_CHAIN (set)) + { + tree selectors = NULL_TREE; + for (tree sel = OMP_TSS_TRAIT_SELECTORS (set); sel; + sel = TREE_CHAIN (sel)) + { + enum omp_ts_code code = OMP_TS_CODE (sel); + tree properties = NULL_TREE; + tree score = OMP_TS_SCORE (sel); + tree t; + + if (score) + { + score = tsubst_expr (score, args, complain, in_decl); + score = fold_non_dependent_expr (score); + if (!INTEGRAL_TYPE_P (TREE_TYPE (score)) + || TREE_CODE (score) != INTEGER_CST) + { + error_at (cp_expr_loc_or_input_loc (score), + "% argument must " + "be constant integer expression"); + score = NULL_TREE; + } + else if (tree_int_cst_sgn (score) < 0) + { + error_at (cp_expr_loc_or_input_loc (score), + "% argument must be non-negative"); + score = NULL_TREE; + } + } + + switch (omp_ts_map[OMP_TS_CODE (sel)].tp_type) + { + case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR: + case OMP_TRAIT_PROPERTY_BOOL_EXPR: + t = tsubst_expr (OMP_TP_VALUE (OMP_TS_PROPERTIES (sel)), + args, complain, in_decl); + t = fold_non_dependent_expr (t); + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) + error_at (cp_expr_loc_or_input_loc (t), + "property must be integer expression"); + properties = make_trait_property (NULL_TREE, t, NULL_TREE); + break; + case OMP_TRAIT_PROPERTY_CLAUSE_LIST: + if (OMP_TS_CODE (sel) == OMP_TRAIT_CONSTRUCT_SIMD) + properties = tsubst_omp_clauses (OMP_TS_PROPERTIES (sel), + C_ORT_OMP_DECLARE_SIMD, + args, complain, in_decl); + break; + default: + /* Nothing to do here, just copy. */ + for (tree prop = OMP_TS_PROPERTIES (sel); + prop; prop = TREE_CHAIN (prop)) + properties = make_trait_property (OMP_TP_NAME (prop), + OMP_TP_VALUE (prop), + properties); + } + selectors = make_trait_selector (code, score, properties, selectors); + } + new_ctx = make_trait_set_selector (OMP_TSS_CODE (set), + nreverse (selectors), + new_ctx); + } + return nreverse (new_ctx); +} + + /* Like tsubst_expr, but unshare TREE_LIST nodes. */ static tree @@ -19406,6 +19479,52 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl) } break; + case OMP_METADIRECTIVE: + { + tree variants = NULL_TREE; + for (tree v = OMP_METADIRECTIVE_VARIANTS (t); v; v = TREE_CHAIN (v)) + { + tree ctx = OMP_METADIRECTIVE_VARIANT_SELECTOR (v); + tree directive = OMP_METADIRECTIVE_VARIANT_DIRECTIVE (v); + tree body = OMP_METADIRECTIVE_VARIANT_BODY (v); + tree s; + + /* CTX is null if this is the default variant. */ + if (ctx) + { + ctx = tsubst_omp_context_selector (ctx, args, complain, + in_decl); + /* Remove the selector from further consideration if it can be + evaluated as a non-match at this point. */ + if (omp_context_selector_matches (ctx, true, true) == 0) + continue; + } + s = push_stmt_list (); + RECUR (directive); + directive = pop_stmt_list (s); + if (body) + { + s = push_stmt_list (); + RECUR (body); + body = pop_stmt_list (s); + } + variants + = chainon (variants, + make_omp_metadirective_variant (ctx, directive, + body)); + } + t = copy_node (t); + OMP_METADIRECTIVE_VARIANTS (t) = variants; + + /* Try to resolve the metadirective early. */ + vec candidates + = omp_early_resolve_metadirective (t); + if (!candidates.is_empty ()) + t = c_omp_expand_metadirective (candidates); + add_stmt (t); + break; + } + case TRANSACTION_EXPR: { int flags = 0; diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C new file mode 100644 index 00000000000..a0b09088d3a --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C @@ -0,0 +1,40 @@ +// { dg-do compile { target c++11 } } + +#define N 100 + +void f (int a[], int b[], int c[]) +{ + int i; + + [[omp::directive (metadirective + default (teams loop) + default (parallel loop))]] /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + [[omp::directive (metadirective + default (bad_directive))]] /* { dg-error "unknown directive name before '\\)' token" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + [[omp::directive (metadirective + default (teams loop) + where (device={arch("nvptx")}: parallel loop))]] /* { dg-error "'where' is not valid for 'metadirective'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + [[omp::directive (metadirective + default (teams loop) + when (device={arch("nvptx")} parallel loop))]] /* { dg-error "expected ':' before 'parallel'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + [[omp::directive (metadirective + default (metadirective default (flush)))]] /* { dg-error "metadirectives cannot be used as variants of a 'metadirective' before 'default'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + /* Test improperly nested metadirectives - even though the second + metadirective resolves to 'omp nothing', that is not the same as there + being literally nothing there. */ + [[omp::directive (metadirective + when (implementation={vendor("gnu")}: parallel for))]] + [[omp::directive (metadirective /* { dg-error "'#pragma' is not allowed here" } */ + when (implementation={vendor("cray")}: parallel for))]] + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; +} diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C new file mode 100644 index 00000000000..44c87df1776 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C @@ -0,0 +1,74 @@ +// { dg-do compile { target c++11 } } + +#define N 100 + +int main (void) +{ + int x = 0; + int y = 0; + + /* Test implicit default (nothing). */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: barrier))]] + x = 1; + + /* Test with multiple standalone directives. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: barrier), + default (flush))]] + x = 1; + + /* Test combining a standalone directive with one that takes a statement + body. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: parallel), + default (barrier))]] + x = 1; + + /* Test combining a standalone directive with one that takes a for loop. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: parallel for), + default (barrier))]] + for (int i = 0; i < N; i++) + x += i; + + /* Test combining a directive that takes a for loop with one that takes + a regular statement body. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: parallel for), + default (parallel))]] + for (int i = 0; i < N; i++) + x += i; + + /* Test labels inside statement body. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: teams num_teams(512)), + when (device={arch("gcn")}: teams num_teams(256)), + default (teams num_teams(4)))]] + { + if (x) + goto l1; + else + goto l2; + l1: ; + l2: ; + } + + /* Test local labels inside statement body. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: teams num_teams(512)), + when (device={arch("gcn")}: teams num_teams(256)), + default (teams num_teams(4)))]] + { + //__label__ l1, l2; + + if (x) + goto l1; + else + goto l2; + l1: ; + l2: ; + } + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C new file mode 100644 index 00000000000..0c2bbdd2f10 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C @@ -0,0 +1,31 @@ +// { dg-do compile { target c++11 } } +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ +/* { dg-additional-options "-fdump-tree-optimized" } */ + +#define N 100 + +void f (int x[], int y[], int z[]) +{ + int i; + + [[omp::sequence (directive (target map(to: x, y) map(from: z)), + directive (metadirective + when (device={arch("nvptx")}: teams loop) + default (parallel loop)))]] + for (i = 0; i < N; i++) + z[i] = x[i] * y[i]; +} + +/* The metadirective should be resolved after Gimplification. */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "when \\(device = .*arch.*nvptx.*\\):" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C new file mode 100644 index 00000000000..43b939be43b --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C @@ -0,0 +1,41 @@ +// { dg-do compile { target c++11 } } + +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 + +#pragma omp declare target +void f(double a[], double x) { + int i; + + [[omp::directive (metadirective + when (construct={target}: distribute parallel for) + default (parallel for simd))]] + for (i = 0; i < N; i++) + a[i] = x * i; +} +#pragma omp end declare target + + int main() +{ + double a[N]; + + #pragma omp target teams map(from: a[0:N]) + f (a, 3.14159); + + /* TODO: This does not execute a version of f with the default clause + active as might be expected. */ + f (a, 2.71828); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ + + return 0; + } + + /* The metadirective should be resolved during Gimplification. */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "when \\(construct = .*target.*\\):" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C new file mode 100644 index 00000000000..1a9cee15be3 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C @@ -0,0 +1,24 @@ +// { dg-do compile { target c++11 } } +/* { dg-additional-options "-fdump-tree-original" } */ + +#define N 100 + +void f (int a[], int flag) +{ + int i; + [[omp::directive (metadirective + when (user={condition(flag)}: + target teams distribute parallel for map(from: a[0:N])) + default (parallel for))]] + for (i = 0; i < N; i++) + a[i] = i; +} + +/* The metadirective should be resolved at parse time. */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp for" 2 "original" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C new file mode 100644 index 00000000000..8a104ff2ebe --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C @@ -0,0 +1,31 @@ +// { dg-do compile { target c++11 } } +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 + +void bar (int a[], int run_parallel, int run_guided) +{ + [[omp::directive (metadirective + when (user={condition(run_parallel)}: parallel))]] + { + int i; + [[omp::directive (metadirective + when (construct={parallel}, user={condition(run_guided)}: + for schedule(guided)) + when (construct={parallel}: for schedule(static)))]] + for (i = 0; i < N; i++) + a[i] = i; + } + } + +/* The outer metadirective should be resolved at parse time. */ +/* The inner metadirective should be resolved during Gimplificiation. */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 2 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp for" 4 "original" } } */ +/* { dg-final { scan-tree-dump-times "when \\(construct = .parallel" 4 "original" } } */ +/* { dg-final { scan-tree-dump-times "default:" 2 "original" } } */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C new file mode 100644 index 00000000000..33861ec077e --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C @@ -0,0 +1,31 @@ +// { dg-do compile { target c++11 } } +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 256 + +void f (int a[], int num) +{ + int i; + + [[omp::directive (metadirective + when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: + target parallel for map(tofrom: a[0:N])) + when (target_device={device_num(num), kind("gpu"), + arch("amdgcn"), isa("gfx906")}: + target parallel for) + when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: + parallel for))]] + for (i = 0; i < N; i++) + a[i] += i; + + [[omp::directive (metadirective + when (target_device={kind("gpu"), arch("nvptx")}: + target parallel for map(tofrom: a[0:N])))]] + for (i = 0; i < N; i++) + a[i] += i; +} + +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"gpu.x00\"\\\[0\\\], &\"amdgcn.x00\"\\\[0\\\], &\"gfx906.x00\"\\\[0\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"gpu.x00\"\\\[0\\\], &\"nvptx.x00\"\\\[0\\\], 0B\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"cpu.x00\"\\\[0\\\], &\"x86_64.x00\"\\\[0\\\], 0B\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(-2, &\"gpu.x00\"\\\[0\\\], &\"nvptx.x00\"\\\[0\\\], 0B\\)" "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C new file mode 100644 index 00000000000..dcb3c365b80 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C @@ -0,0 +1,16 @@ +// { dg-do compile { target c++11 } } + +#define N 256 + +void f () +{ + int i; + int a[N]; + + [[omp::directive (metadirective + when( device={kind(nohost)}: nothing ) + when( device={arch("nvptx")}: nothing) + default( parallel for))]] + for (i = 0; i < N; i++) + a[i] = i; +} diff --git a/libgomp/testsuite/libgomp.c++/metadirective-template-1.C b/libgomp/testsuite/libgomp.c++/metadirective-template-1.C new file mode 100644 index 00000000000..2d826574861 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/metadirective-template-1.C @@ -0,0 +1,37 @@ +# include + +template int +fib (int n) +{ + int i, j; + if (n < 2) + return n; + else if ( tasking && n < 8 ) // serial/taskless cutoff for n<8 + return fib (n); + else + { +#pragma omp metadirective \ + when (user = {condition (tasking)}: task shared(i)) + i = fib (n - 1); +#pragma omp metadirective \ + when (user = {condition (score(10): tasking)}: task shared(j)) + j = fib (n - 2); +#pragma omp metadirective \ + when (user = {condition (tasking)}: taskwait) + return i + j; + } +} + +int main () +{ + int n = 15, o = 610; +#pragma omp parallel +#pragma omp single + { + if (fib (n) != o) + __builtin_abort (); + if (fib (n) != o) + __builtin_abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/metadirective-template-2.C b/libgomp/testsuite/libgomp.c++/metadirective-template-2.C new file mode 100644 index 00000000000..9134fefc7ac --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/metadirective-template-2.C @@ -0,0 +1,41 @@ +# include + +template int +fib (int n, bool flag) +{ + int i, j; + if (n < 2) + return n; + else if ( tasking && flag && n < 8 ) // serial/taskless cutoff for n<8 + return fib (n, false); + else + { +#pragma omp metadirective \ + when (user = {condition (tasking && flag)}: task shared(i)) + i = fib (n - 1, flag); +#pragma omp metadirective \ + when (user = {condition (score(10): tasking && flag)}: task shared(j)) + j = fib (n - 2, flag); +#pragma omp metadirective \ + when (user = {condition (tasking && flag)}: taskwait) + return i + j; + } +} + +int main () +{ + int n = 15, o = 610; +#pragma omp parallel +#pragma omp single + { + if (fib (n, true) != o) + __builtin_abort (); + if (fib (n, false) != o) + __builtin_abort (); + if (fib (n, false) != o) + __builtin_abort (); + if (fib (n, true) != o) + __builtin_abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/metadirective-template-3.C b/libgomp/testsuite/libgomp.c++/metadirective-template-3.C new file mode 100644 index 00000000000..cc48fde0fda --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/metadirective-template-3.C @@ -0,0 +1,41 @@ +# include + +template int +fib (int n, bool flag) +{ + int i, j; + if (n < 2) + return n; + else if ( tasking && flag && n < 8 ) // serial/taskless cutoff for n<8 + return fib (n, false); + else + { +#pragma omp metadirective \ + when (user = {condition (tasking && flag)}: task shared(i)) \ + when (user = {condition (!tasking && !flag)}: nothing) \ + otherwise (error at(execution) message("oops 1")) + i = fib (n - 1, flag); +#pragma omp metadirective \ + when (user = {condition (score(10): tasking && flag)}: task shared(j)) \ + when (user = {condition (tasking || flag)} : \ + error at(execution) message ("oops 2")) + j = fib (n - 2, flag); +#pragma omp metadirective \ + when (user = {condition (tasking && flag)}: taskwait) + return i + j; + } +} + +int main () +{ + int n = 15, o = 610; +#pragma omp parallel +#pragma omp single + { + if (fib (n, true) != o) + __builtin_abort (); + if (fib (n, false) != o) + __builtin_abort (); + } + return 0; +}