From patchwork Fri Jul 12 14:11:51 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 1959901 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=JpyZzQvg; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; 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 [8.43.85.97]) (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 4WLDCy0qsFz1xpd for ; Sat, 13 Jul 2024 00:16:14 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 4E03138323F5 for ; Fri, 12 Jul 2024 14:16:12 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lf1-x130.google.com (mail-lf1-x130.google.com [IPv6:2a00:1450:4864:20::130]) by sourceware.org (Postfix) with ESMTPS id CA128385DDE9 for ; Fri, 12 Jul 2024 14:12:58 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CA128385DDE9 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 CA128385DDE9 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::130 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; cv=none; b=ICwxXSBZ7pNxA7CbCwdz0f2fn6j7MSx8Hepz76k1282X8gF8E618VkxGCaWyLPe2ZSsbfIek95aYjiURBSsH/DFxJXWi2zvWe6la53AP486HozmiOFruwUwOmpvItmxY4kKf/UJvNwiTPiJ/59VKFgBk0olxo+GGz6306O1yriI= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1720793583; c=relaxed/simple; bh=MQnBsEmyR+kix6dYpQL77iF63/+MnECZUhna83nGtzA=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=l2aD8hffTagfW5E0mCMSZIAHRR75aREVvE8N4vKOY1mSmPaIVcBU4SwDwkk0RDy83TmzkD4RLNC0WUwqSCi5VcVDg4nx31vKq5m8t3BR/f2/ReJYDNS3FPt9dJiUkDMMCp+8phhU8eRMwsgogCU7Qn3Yv522XZRPDo0y9G36e/0= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lf1-x130.google.com with SMTP id 2adb3069b0e04-52e99060b41so2189260e87.2 for ; Fri, 12 Jul 2024 07:12:58 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1720793576; x=1721398376; 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=tLNfdHs718tFAuvaMDy34TpGxYu1pgZ0ZRdWrjF4aX8=; b=JpyZzQvgrXCt11s/AhFJuYViody/XzlRbte+UrdQCOErWZqb9Rthf/N04K2xv6zH16 jLKLe0in74Lto9CvEakcS9yFU7rKynmFsG1J+qc1k9VrxLtr0YNNKxsvbvZbbvHBUFmT 26h0rNoiJ3U3RnntpfGJUl0vLvVK+0XOliO5oDoCEAyutxTQOVlI/E6T3eBqT9P5QPCb 7iGEVtqHVZjo2j8Bp2/Sw44wAMR6UAZWHxhh/YDbD2YX3AFbxfGESVAclQRRgdbBQkdY 72R4Vogiq9bnuOBF2doEB+z+2ySGXltGx6NVWv09pexSR4D4hSSNJ59DpN7i6wUU0RP1 VTIw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1720793576; x=1721398376; 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=tLNfdHs718tFAuvaMDy34TpGxYu1pgZ0ZRdWrjF4aX8=; b=A4R2lzJ8DSde0ol8RRF5AmoRFeMjmHXEnETyXaKcOS1Xnp+Hn34SvFmMK0tCJsIWZN rsG+OKK6Tb6wEX2EbZqeiaM9R9uTaR+6a5V9vJVsRhv2H5Lmly6DRKlMTQPrSE1DzpxN deLmfYzB4hGt493BQGwxMOKQ6F4LDukE71fmrW60q9hTjq3kdnk6Rtc1jq+rNsngO7rT Pr7V16SxuisDquaB8URQ3TqRgVCUH4IK+ovVdNMnDEfAtxcix+wnAiRCdGlJu53mnfUZ 4SlWdI0c+X9bDux+5WS28n00TEIyCmNdftTs2ggqQxqvNOgtL/ZpOAFfOhkeU2ZJ7Ceu WbGQ== X-Gm-Message-State: AOJu0YwWpEf1YHZgf51f/j1b+6PcJCG0ofrkLo0zrca6RqDn9gsf6ako NFgBOuuUHvKcaxo40+q099Ii5xFI+lvA09FKPvDXeDoi10iZfqhWg4D3rmRZOxvcEbuKolazGsV z X-Google-Smtp-Source: AGHT+IE3TzwEmq9W9wvwF0OS1DBWTgvh2ZW3LbYMi3j7qUS5Rjr9e6PtGmcTROGIl8qYfmq+LSlPTg== X-Received: by 2002:a05:6512:3b84:b0:52e:9ecd:3465 with SMTP id 2adb3069b0e04-52eb99d4e22mr8299286e87.57.1720793575992; Fri, 12 Jul 2024 07:12:55 -0700 (PDT) Received: from localhost.localdomain ([169.155.255.128]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-4279f25b946sm24680325e9.19.2024.07.12.07.12.55 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 Jul 2024 07:12:55 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v2 4/8] OpenMP: C front-end support for dispatch + adjust_args Date: Fri, 12 Jul 2024 16:11:51 +0200 Message-ID: <20240712141155.255186-5-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20240712141155.255186-1-parras@baylibre.com> References: <20240712141155.255186-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.9 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 support to the C front-end to parse the `dispatch` construct and the `adjust_args` clause. It also includes some common C/C++ bits for pragmas and attributes. Additional common C/C++ testcases are in a later patch in the series. gcc/c-family/ChangeLog: * c-attribs.cc (c_common_gnu_attributes): Add attribute for adjust_args need_device_ptr. * c-omp.cc (c_omp_directives): Uncomment dispatch. * c-pragma.cc (omp_pragmas): Add dispatch. * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_DISPATCH. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_NOCONTEXT and PRAGMA_OMP_CLAUSE_NOVARIANTS. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_dispatch): New function. (c_parser_omp_clause_name): Handle nocontext and novariants clauses. (c_parser_omp_clause_novariants): New function. (c_parser_omp_clause_nocontext): Likewise. (c_parser_omp_all_clauses): Handle nocontext and novariants clauses. (c_parser_omp_dispatch_body): New function adapted from c_parser_expr_no_commas. (OMP_DISPATCH_CLAUSE_MASK): Define. (c_parser_omp_dispatch): New function. (c_finish_omp_declare_variant): Parse adjust_args. (c_parser_omp_construct): Handle PRAGMA_OMP_DISPATCH. * c-typeck.cc (c_finish_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. gcc/testsuite/ChangeLog: * gcc.dg/gomp/adjust-args-1.c: New test. * gcc.dg/gomp/dispatch-1.c: New test. --- gcc/c-family/c-attribs.cc | 2 + gcc/c-family/c-omp.cc | 4 +- gcc/c-family/c-pragma.cc | 1 + gcc/c-family/c-pragma.h | 3 + gcc/c/c-parser.cc | 496 +++++++++++++++++++--- gcc/c/c-typeck.cc | 2 + gcc/testsuite/gcc.dg/gomp/adjust-args-1.c | 32 ++ gcc/testsuite/gcc.dg/gomp/dispatch-1.c | 53 +++ libgomp/testsuite/libgomp.c/dispatch-1.c | 76 ++++ 9 files changed, 609 insertions(+), 60 deletions(-) create mode 100644 gcc/testsuite/gcc.dg/gomp/adjust-args-1.c create mode 100644 gcc/testsuite/gcc.dg/gomp/dispatch-1.c create mode 100644 libgomp/testsuite/libgomp.c/dispatch-1.c diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc index f9b229aba7f..1cb49d7b911 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -560,6 +560,8 @@ const struct attribute_spec c_common_gnu_attributes[] = handle_omp_declare_variant_attribute, NULL }, { "omp declare variant variant", 0, -1, true, false, false, false, handle_omp_declare_variant_attribute, NULL }, + { "omp declare variant adjust_args need_device_ptr", 0, -1, true, false, false, false, + handle_omp_declare_variant_attribute, NULL }, { "simd", 0, 1, true, false, false, false, handle_simd_attribute, NULL }, { "omp declare target", 0, -1, true, false, false, false, diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index b5ce1466e5d..c74a9fb2691 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4299,8 +4299,8 @@ const struct c_omp_directive c_omp_directives[] = { C_OMP_DIR_DECLARATIVE, false }, { "depobj", nullptr, nullptr, PRAGMA_OMP_DEPOBJ, C_OMP_DIR_STANDALONE, false }, - /* { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH, - C_OMP_DIR_CONSTRUCT, false }, */ + { "dispatch", nullptr, nullptr, PRAGMA_OMP_DISPATCH, + C_OMP_DIR_DECLARATIVE, false }, { "distribute", nullptr, nullptr, PRAGMA_OMP_DISTRIBUTE, C_OMP_DIR_CONSTRUCT, true }, { "end", "assumes", nullptr, PRAGMA_OMP_END, diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc index 25251c2b69f..b956819c0a5 100644 --- a/gcc/c-family/c-pragma.cc +++ b/gcc/c-family/c-pragma.cc @@ -1526,6 +1526,7 @@ static const struct omp_pragma_def omp_pragmas[] = { { "cancellation", PRAGMA_OMP_CANCELLATION_POINT }, { "critical", PRAGMA_OMP_CRITICAL }, { "depobj", PRAGMA_OMP_DEPOBJ }, + { "dispatch", PRAGMA_OMP_DISPATCH }, { "error", PRAGMA_OMP_ERROR }, { "end", PRAGMA_OMP_END }, { "flush", PRAGMA_OMP_FLUSH }, diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 2ebde06c471..6b6826b2426 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -55,6 +55,7 @@ enum pragma_kind { PRAGMA_OMP_CRITICAL, PRAGMA_OMP_DECLARE, PRAGMA_OMP_DEPOBJ, + PRAGMA_OMP_DISPATCH, PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_ERROR, PRAGMA_OMP_END, @@ -135,9 +136,11 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_LINK, PRAGMA_OMP_CLAUSE_MAP, PRAGMA_OMP_CLAUSE_MERGEABLE, + PRAGMA_OMP_CLAUSE_NOCONTEXT, PRAGMA_OMP_CLAUSE_NOGROUP, PRAGMA_OMP_CLAUSE_NONTEMPORAL, PRAGMA_OMP_CLAUSE_NOTINBRANCH, + PRAGMA_OMP_CLAUSE_NOVARIANTS, PRAGMA_OMP_CLAUSE_NOWAIT, PRAGMA_OMP_CLAUSE_NUM_TASKS, PRAGMA_OMP_CLAUSE_NUM_TEAMS, diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 12c5ed5d92c..8a6653057ef 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1740,6 +1740,8 @@ static void c_parser_omp_assumption_clauses (c_parser *, bool); static void c_parser_omp_allocate (c_parser *); static void c_parser_omp_assumes (c_parser *); static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *); +static tree +c_parser_omp_dispatch (location_t, c_parser *); static void c_parser_oacc_routine (c_parser *, enum pragma_context); /* These Objective-C parser functions are only ever called when @@ -15044,6 +15046,8 @@ c_parser_omp_clause_name (c_parser *parser) case 'n': if (!strcmp ("no_create", p)) result = PRAGMA_OACC_CLAUSE_NO_CREATE; + else if (!strcmp ("nocontext", p)) + result = PRAGMA_OMP_CLAUSE_NOCONTEXT; else if (!strcmp ("nogroup", p)) result = PRAGMA_OMP_CLAUSE_NOGROUP; else if (!strcmp ("nohost", p)) @@ -15052,6 +15056,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_NONTEMPORAL; else if (!strcmp ("notinbranch", p)) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; + else if (!strcmp ("novariants", p)) + result = PRAGMA_OMP_CLAUSE_NOVARIANTS; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; else if (!strcmp ("num_gangs", p)) @@ -19319,6 +19325,60 @@ c_parser_omp_clause_partial (c_parser *parser, tree list) return c; } +/* OpenMP 5.1 + novariants ( scalar-expression ) */ + +static tree +c_parser_omp_clause_novariants (c_parser *parser, tree list) +{ + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + location_t loc = c_parser_peek_token (parser)->location; + c_expr expr = c_parser_expr_no_commas (parser, NULL); + tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value; + t = c_objc_common_truthvalue_conversion (loc, t); + t = c_fully_fold (t, false, NULL); + parens.skip_until_found_close (parser); + + check_no_duplicate_clause (list, OMP_CLAUSE_NOVARIANTS, "novariants"); + + tree c = build_omp_clause (loc, OMP_CLAUSE_NOVARIANTS); + OMP_CLAUSE_NOVARIANTS_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + + return list; +} + +/* OpenMP 5.1 + nocontext ( scalar-expression ) */ + +static tree +c_parser_omp_clause_nocontext (c_parser *parser, tree list) +{ + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + location_t loc = c_parser_peek_token (parser)->location; + c_expr expr = c_parser_expr_no_commas (parser, NULL); + tree t = convert_lvalue_to_rvalue (loc, expr, true, true).value; + t = c_objc_common_truthvalue_conversion (loc, t); + t = c_fully_fold (t, false, NULL); + parens.skip_until_found_close (parser); + + check_no_duplicate_clause (list, OMP_CLAUSE_NOCONTEXT, "nocontext"); + + tree c = build_omp_clause (loc, OMP_CLAUSE_NOCONTEXT); + OMP_CLAUSE_NOCONTEXT_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + + return list; +} + /* OpenMP 5.0: detach ( event-handle ) */ @@ -19938,6 +19998,14 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "partial"; clauses = c_parser_omp_clause_partial (parser, clauses); break; + case PRAGMA_OMP_CLAUSE_NOVARIANTS: + c_name = "novariants"; + clauses = c_parser_omp_clause_novariants (parser, clauses); + break; + case PRAGMA_OMP_CLAUSE_NOCONTEXT: + c_name = "nocontext"; + clauses = c_parser_omp_clause_nocontext (parser, clauses); + break; default: c_parser_error (parser, "expected an OpenMP clause"); goto saw_error; @@ -23748,6 +23816,168 @@ c_parser_omp_scope (location_t loc, c_parser *parser, bool *if_p) return add_stmt (stmt); } +// Adapted from c_parser_expr_no_commas +static tree +c_parser_omp_dispatch_body (c_parser *parser) +{ + struct c_expr lhs, rhs, ret; + struct c_expr orig_expr; + location_t expr_loc = c_parser_peek_token (parser)->location; + source_range tok_range = c_parser_peek_token (parser)->get_range (); + location_t sizeof_arg_loc[3]; + tree sizeof_arg[3]; + vec *exprlist; + vec arg_loc = vNULL; + vec *origtypes = NULL; + unsigned int literal_zero_mask; + location_t start; + location_t finish; + + lhs = c_parser_conditional_expression (parser, NULL, NULL); + if (TREE_CODE (lhs.value) == CALL_EXPR) + return lhs.value; + else + { + location_t op_location = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_EQ, "expected %<=%>")) + return error_mark_node; + + /* Parse function name*/ + if (!c_parser_next_token_is (parser, CPP_NAME)) + { + c_parser_error (parser, "expected a function name"); + rhs.set_error (); + return rhs.value; + } + expr_loc = c_parser_peek_token (parser)->location; + tree id = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return error_mark_node; + + rhs.value = build_external_ref (expr_loc, id, true, &rhs.original_type); + set_c_expr_source_range (&rhs, tok_range); + /* Parse argument list */ + { + for (int i = 0; i < 3; i++) + { + sizeof_arg[i] = NULL_TREE; + sizeof_arg_loc[i] = UNKNOWN_LOCATION; + } + literal_zero_mask = 0; + if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN)) + exprlist = NULL; + else + exprlist = c_parser_expr_list (parser, true, false, &origtypes, + sizeof_arg_loc, sizeof_arg, &arg_loc, + &literal_zero_mask); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + } + orig_expr = rhs; + mark_exp_read (rhs.value); + if (warn_sizeof_pointer_memaccess) + sizeof_pointer_memaccess_warning (sizeof_arg_loc, rhs.value, exprlist, + sizeof_arg, + sizeof_ptr_memacc_comptypes); + if (TREE_CODE (rhs.value) == FUNCTION_DECL) + { + if (fndecl_built_in_p (rhs.value, BUILT_IN_MEMSET) + && vec_safe_length (exprlist) == 3) + { + tree arg0 = (*exprlist)[0]; + tree arg2 = (*exprlist)[2]; + warn_for_memset (expr_loc, arg0, arg2, literal_zero_mask); + } + if (warn_absolute_value + && fndecl_built_in_p (rhs.value, BUILT_IN_NORMAL) + && vec_safe_length (exprlist) == 1) + warn_for_abs (expr_loc, rhs.value, (*exprlist)[0]); + if (parser->omp_for_parse_state + && parser->omp_for_parse_state->in_intervening_code + && omp_runtime_api_call (rhs.value)) + { + error_at (expr_loc, "calls to the OpenMP runtime API are " + "not permitted in intervening code"); + parser->omp_for_parse_state->fail = true; + } + } + + start = rhs.get_start (); + finish = parser->tokens_buf[0].get_finish (); + rhs.value = c_build_function_call_vec (expr_loc, arg_loc, rhs.value, + exprlist, origtypes); + set_c_expr_source_range (&rhs, start, finish); + rhs.m_decimal = 0; + + rhs.original_code = ERROR_MARK; + if (TREE_CODE (rhs.value) == INTEGER_CST + && TREE_CODE (orig_expr.value) == FUNCTION_DECL + && fndecl_built_in_p (orig_expr.value, BUILT_IN_CONSTANT_P)) + rhs.original_code = C_MAYBE_CONST_EXPR; + rhs.original_type = NULL; + if (exprlist) + { + release_tree_vector (exprlist); + release_tree_vector (origtypes); + } + arg_loc.release (); + + /* Build assignment */ + rhs = convert_lvalue_to_rvalue (expr_loc, rhs, true, true); + ret.value + = build_modify_expr (op_location, lhs.value, lhs.original_type, + NOP_EXPR, expr_loc, rhs.value, rhs.original_type); + ret.m_decimal = 0; + set_c_expr_source_range (&ret, lhs.get_start (), rhs.get_finish ()); + ret.original_code = MODIFY_EXPR; + ret.original_type = NULL; + return ret.value; + } +} + +/* OpenMP 5.1: + # pragma omp dispatch dispatch-clause[optseq] new-line + expression-stmt + + LOC is the location of the #pragma. +*/ + +#define OMP_DISPATCH_CLAUSE_MASK \ + ((OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) + +static tree +c_parser_omp_dispatch (location_t loc, c_parser *parser) +{ + tree stmt = make_node (OMP_DISPATCH); + SET_EXPR_LOCATION (stmt, loc); + TREE_TYPE (stmt) = void_type_node; + + OMP_DISPATCH_CLAUSES (stmt) + = c_parser_omp_all_clauses (parser, OMP_DISPATCH_CLAUSE_MASK, + "#pragma omp dispatch"); + + // Parse body as expression statement + loc = c_parser_peek_token (parser)->location; + tree dispatch_body = c_parser_omp_dispatch_body (parser); + if (dispatch_body == error_mark_node) + { + inform (loc, "%<#pragma omp dispatch%> must be followed by a function " + "call with optional assignment"); + c_parser_skip_to_end_of_block_or_statement (parser); + return NULL_TREE; + } + + c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>"); + OMP_DISPATCH_BODY (stmt) = dispatch_body; + + return add_stmt (stmt); +} + /* OpenMP 3.0: # pragma omp task task-clause[optseq] new-line @@ -24728,6 +24958,10 @@ check_clauses: OpenMP 5.0: # pragma omp declare variant (identifier) match(context-selector) new-line + + OpenMP 5.1 + # pragma omp declare variant (identifier) match(context-selector) \ + adjust_args(adjust-op:argument-list) new-line */ #define OMP_DECLARE_SIMD_CLAUSE_MASK \ @@ -25191,77 +25425,217 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms) parens.require_close (parser); - if (c_parser_next_token_is (parser, CPP_COMMA) - && c_parser_peek_2nd_token (parser)->type == CPP_NAME) - c_parser_consume_token (parser); + vec adjust_args_list = vNULL; + bool has_match = false, has_adjust_args = false; + location_t adjust_args_loc = UNKNOWN_LOCATION; - const char *clause = ""; - location_t match_loc = c_parser_peek_token (parser)->location; - if (c_parser_next_token_is (parser, CPP_NAME)) - clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); - if (strcmp (clause, "match")) + do { - c_parser_error (parser, "expected %"); - goto fail; - } + if (c_parser_next_token_is (parser, CPP_COMMA) + && c_parser_peek_2nd_token (parser)->type == CPP_NAME) + c_parser_consume_token (parser); - c_parser_consume_token (parser); + const char *clause = ""; + location_t match_loc = c_parser_peek_token (parser)->location; + if (c_parser_next_token_is (parser, CPP_NAME)) + clause = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); - if (!parens.require_open (parser)) - goto fail; + enum clause + { + match, + adjust_args + } ccode; - if (parms == NULL_TREE) - parms = error_mark_node; - - tree ctx = c_parser_omp_context_selector_specification (parser, parms); - if (ctx == error_mark_node) - goto fail; - ctx = omp_check_context_selector (match_loc, ctx); - if (ctx != error_mark_node && variant != error_mark_node) - { - if (TREE_CODE (variant) != FUNCTION_DECL) + if (strcmp (clause, "match") == 0) + ccode = match; + else if (strcmp (clause, "adjust_args") == 0) { - error_at (token->location, "variant %qD is not a function", variant); - variant = error_mark_node; + ccode = adjust_args; + adjust_args_loc = match_loc; } - else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, - OMP_TRAIT_CONSTRUCT_SIMD) - && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant))) + else { - error_at (token->location, "variant %qD and base %qD have " - "incompatible types", variant, fndecl); - variant = error_mark_node; + c_parser_error (parser, "expected % or %"); + goto fail; } - else if (fndecl_built_in_p (variant) - && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), - "__builtin_", strlen ("__builtin_")) == 0 - || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), - "__sync_", strlen ("__sync_")) == 0 - || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), - "__atomic_", strlen ("__atomic_")) == 0)) + + c_parser_consume_token (parser); + + if (!parens.require_open (parser)) + goto fail; + + if (parms == NULL_TREE) + parms = error_mark_node; + + if (ccode == match) { - error_at (token->location, "variant %qD is a built-in", variant); - variant = error_mark_node; - } - if (variant != error_mark_node) - { - C_DECL_USED (variant) = 1; - tree construct - = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT); - omp_mark_declare_variant (match_loc, variant, construct); - if (omp_context_selector_matches (ctx)) + has_match = true; + tree ctx + = c_parser_omp_context_selector_specification (parser, parms); + if (ctx == error_mark_node) + goto fail; + ctx = omp_check_context_selector (match_loc, ctx); + if (ctx != error_mark_node && variant != error_mark_node) { - tree attr - = tree_cons (get_identifier ("omp declare variant base"), - build_tree_list (variant, ctx), - DECL_ATTRIBUTES (fndecl)); - DECL_ATTRIBUTES (fndecl) = attr; + if (TREE_CODE (variant) != FUNCTION_DECL) + { + error_at (token->location, "variant %qD is not a function", + variant); + variant = error_mark_node; + } + else if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_SIMD) + && !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant))) + { + error_at (token->location, + "variant %qD and base %qD have " + "incompatible types", + variant, fndecl); + variant = error_mark_node; + } + else if (fndecl_built_in_p (variant) + && (strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), + "__builtin_", strlen ("__builtin_")) + == 0 + || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), + "__sync_", strlen ("__sync_")) + == 0 + || strncmp (IDENTIFIER_POINTER (DECL_NAME (variant)), + "__atomic_", strlen ("__atomic_")) + == 0)) + { + error_at (token->location, "variant %qD is a built-in", + variant); + variant = error_mark_node; + } + if (variant != error_mark_node) + { + C_DECL_USED (variant) = 1; + tree construct + = omp_get_context_selector_list (ctx, + OMP_TRAIT_SET_CONSTRUCT); + omp_mark_declare_variant (match_loc, variant, construct); + if (omp_context_selector_matches (ctx)) + { + tree attr = tree_cons (get_identifier ( + "omp declare variant base"), + build_tree_list (variant, ctx), + DECL_ATTRIBUTES (fndecl)); + DECL_ATTRIBUTES (fndecl) = attr; + } + } } } - } + else if (ccode == adjust_args) + { + has_adjust_args = true; + if (c_parser_next_token_is (parser, CPP_NAME) + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + { + const char *p + = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + if (strcmp (p, "need_device_ptr") == 0 + || strcmp (p, "nothing") == 0) + { + c_parser_consume_token (parser); // need_device_ptr + c_parser_consume_token (parser); // : - parens.require_close (parser); + location_t loc = c_parser_peek_token (parser)->location; + tree list + = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_ERROR, + NULL_TREE); + + tree arg; + for (tree c = list; c != NULL_TREE; c = TREE_CHAIN (c)) + { + tree decl = TREE_PURPOSE (c); + int idx; + for (arg = parms, idx = 0; arg != NULL; + arg = TREE_CHAIN (arg), idx++) + if (arg == decl) + break; + if (arg == NULL_TREE) + { + error_at (loc, "%qD is not a function argument", + decl); + goto fail; + } + if (adjust_args_list.contains (arg)) + { + error_at (loc, "%qD is specified more than once", + decl); + goto fail; + } + if (strcmp (p, "need_device_ptr") == 0 + && TREE_CODE (TREE_TYPE (arg)) != POINTER_TYPE) + { + error_at (loc, "%qD is not a C pointer", decl); + goto fail; + } + adjust_args_list.safe_push (arg); + if (strcmp (p, "need_device_ptr") == 0) + { + tree variant_decl = (TREE_CODE (variant) == NOP_EXPR) + ? TREE_OPERAND (variant, 0) + : variant; + tree variant_parm = DECL_ARGUMENTS (variant_decl); + for (int i = 0; i < idx; i++) + { + variant_parm = TREE_CHAIN (variant_parm); + gcc_assert (variant_parm != NULL); + } + tree attr = tree_cons ( + get_identifier ("omp declare variant adjust_args " + "need_device_ptr"), + NULL_TREE, DECL_ATTRIBUTES (variant_parm)); + DECL_ATTRIBUTES (variant_parm) = attr; + } + } + } + else + { + error_at (c_parser_peek_token (parser)->location, + "expected % or %"); + goto fail; + } + } + else + { + error_at (c_parser_peek_token (parser)->location, + "expected % or % " + "followed by %<:%>"); + goto fail; + } + } + + parens.require_close (parser); + } while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)); c_parser_skip_to_pragma_eol (parser); + + if (has_adjust_args) + { + if (!has_match) + { + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + else + { + tree attr = lookup_attribute ("omp declare variant base", + DECL_ATTRIBUTES (fndecl)); + tree ctx = TREE_VALUE (TREE_VALUE (attr)); + if (!omp_get_context_selector (ctx, OMP_TRAIT_SET_CONSTRUCT, + OMP_TRAIT_CONSTRUCT_DISPATCH)) + error_at ( + adjust_args_loc, + "an % clause can only be specified if the " + "% selector of the construct selector set appears " + "in the % clause"); + } + } } /* Finalize #pragma omp declare simd or #pragma omp declare variant @@ -26078,7 +26452,6 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context) types.release (); } - /* OpenMP 4.0 #pragma omp declare simd declare-simd-clauses[optseq] new-line #pragma omp declare reduction (reduction-id : typename-list : expression) \ @@ -26086,7 +26459,11 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context) #pragma omp declare target new-line OpenMP 5.0 - #pragma omp declare variant (identifier) match (context-selector) */ + #pragma omp declare variant (identifier) match (context-selector) + + OpenMP 5.1 + #pragma omp declare variant (identifier) match (context-selector) \ + adjust_args(adjust-op:argument-list) */ static bool c_parser_omp_declare (c_parser *parser, enum pragma_context context) @@ -27001,6 +27378,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) case PRAGMA_OMP_UNROLL: stmt = c_parser_omp_unroll (loc, parser, if_p); break; + case PRAGMA_OMP_DISPATCH: + stmt = c_parser_omp_dispatch (loc, parser); + break; default: gcc_unreachable (); } diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 36f88fcd03d..5745d0cca47 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -16305,6 +16305,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_INDIRECT: + case OMP_CLAUSE_NOVARIANTS: + case OMP_CLAUSE_NOCONTEXT: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c new file mode 100644 index 00000000000..393a44de8e0 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-1.c @@ -0,0 +1,32 @@ +/* Test parsing of OMP clause adjust_args */ +/* { dg-do compile } */ + +int b; + +int f0 (void *a); +int g (void *a); +int f1 (int); + +#pragma omp declare variant (f0) match (construct={target}) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f2 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (other: a) /* { dg-error "expected 'nothing' or 'need_device_ptr'" } */ +int f3 (int a); +#pragma omp declare variant (f0) adjust_args (nothing: a) /* { dg-error "an 'adjust_args' clause can only be specified if the 'dispatch' selector of the construct selector set appears in the 'match' clause" } */ +int f4 (void *a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args () /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f5 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing) /* { dg-error "expected 'nothing' or 'need_device_ptr' followed by ':'" } */ +int f6 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing:) /* { dg-error "expected identifier before '\\)' token" } */ +int f7 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: z) /* { dg-error "'z' undeclared here \\(not in a function\\)" } */ +int f8 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (need_device_ptr: a) /* { dg-error "'a' is not a C pointer" } */ +int f9 (int a); +#pragma omp declare variant (f1) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (nothing: a) /* { dg-error "'a' is specified more than once" } */ +int f10 (int a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (nothing: a) adjust_args (need_device_ptr: a) /* { dg-error "'a' is specified more than once" } */ +int f11 (void *a); +#pragma omp declare variant (g) match (construct={dispatch}) adjust_args (need_device_ptr: b) /* { dg-error "'b' is not a function argument" } */ +int f12 (void *a); + diff --git a/gcc/testsuite/gcc.dg/gomp/dispatch-1.c b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c new file mode 100644 index 00000000000..c8f45c12be6 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/dispatch-1.c @@ -0,0 +1,53 @@ +/* Test parsing of #pragma omp dispatch */ +/* { dg-do compile } */ + +int f0 (int); + +void f1 (void) +{ + int a, b; + double x; + struct {int a; int b;} s; + int arr[1]; + +#pragma omp dispatch + int c = f0 (a); /* { dg-error "expected expression before 'int'" } */ +#pragma omp dispatch + int f2 (int d); /* { dg-error "expected expression before 'int'" } */ +#pragma omp dispatch + a = b; /* { dg-error "expected '\\(' before ';' token" } */ +#pragma omp dispatch + s.a = f0(a) + b; /* { dg-error "expected ';' before '\\+' token" } */ +#pragma omp dispatch + b = !f0(a); /* { dg-error "expected a function name before '!' token" } */ +#pragma omp dispatch + s.b += f0(s.a); /* { dg-error "expected '=' before '\\+=' token" } */ +#pragma omp dispatch +#pragma omp threadprivate(a) /* { dg-error "expected expression before '#pragma'" } */ + a = f0(b); + +#pragma omp dispatch nocontext(s) /* { dg-error "used struct type value where scalar is required" } */ + f0(a); +#pragma omp dispatch nocontext(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch nocontext(a) nocontext(b) /* { dg-error "too many 'nocontext' clauses" } */ + f0(a); +#pragma omp dispatch novariants(s) /* { dg-error "used struct type value where scalar is required" } */ + f0(a); +#pragma omp dispatch novariants(a, b) /* { dg-error "expected '\\)' before ','" } */ + f0(a); +#pragma omp dispatch novariants(a) novariants(b) /* { dg-error "too many 'novariants' clauses" } */ + f0(a); +#pragma omp dispatch nowait nowait /* { dg-error "too many 'nowait' clauses" } */ + f0(a); +#pragma omp dispatch device(x) /* { dg-error "expected integer expression before end of line" } */ + f0(a); +#pragma omp dispatch device(arr) /* { dg-error "expected integer expression before end of line" } */ + f0(a); +#pragma omp dispatch is_device_ptr(x) /* { dg-error "'is_device_ptr' variable is neither a pointer nor an array" } */ + f0(a); +#pragma omp dispatch is_device_ptr(&x) /* { dg-error "expected identifier before '&' token" } */ + f0(a); +#pragma omp dispatch depend(inout: f0) /* { dg-error "'f0' is not lvalue expression nor array section in 'depend' clause" } */ + f0(a); +} diff --git a/libgomp/testsuite/libgomp.c/dispatch-1.c b/libgomp/testsuite/libgomp.c/dispatch-1.c new file mode 100644 index 00000000000..0efc075a859 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/dispatch-1.c @@ -0,0 +1,76 @@ +// Adapted from OpenMP examples + +#include +#include +#include + +int baz (double *d_bv, const double *d_av, int n) +{ +#pragma omp distribute parallel for + for (int i = 0; i < n; i++) + d_bv[i] = d_av[i] * i; + return -3; +} + +int bar (double *d_bv, const double *d_av, int n) +{ +#pragma omp target is_device_ptr(d_bv, d_av) + for (int i = 0; i < n; i++) + d_bv[i] = d_av[i] * i; + return -2; +} + +#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: bv, av) +#pragma omp declare variant(baz) match(implementation={vendor(gnu)}) +int foo (double *bv, const double *av, int n) +{ + for (int i = 0; i < n; i++) + bv[i] = av[i] * i; + return -1; +} + +int test (int n) +{ + const double e = 2.71828; + + double *av = (double *) malloc (n * sizeof (*av)); + double *bv = (double *) malloc (n * sizeof (*bv)); + double *d_bv = (double *) malloc (n * sizeof (*d_bv)); + + for (int i = 0; i < n; i++) + { + av[i] = e * i; + bv[i] = 0.0; + d_bv[i] = 0.0; + } + + int f, last_dev = omp_get_num_devices () - 1; +#pragma omp target data map(to: av[:n]) map(from: d_bv[:n]) device(last_dev) if (n == 1024) + { + #pragma omp dispatch nocontext(n > 1024) novariants(n < 1024) device(last_dev) + f = foo (d_bv, av, n); + } + + foo (bv, av, n); + for (int i = 0; i < n; i++) + { + if (d_bv[i] != bv[i]) + { + fprintf (stderr, "ERROR at %d: %lf (act) != %lf (exp)\n", i, d_bv[i], bv[i]); + return 1; + } + } + return f; +} + +int +main (void) +{ + int ret = test(1023); + if (ret != -1) return 1; + ret = test(1024); + if (ret != -2) return 1; + ret = test(1025); + if (ret != -3) return 1; + return 0; +}