From patchwork Wed Oct 2 16:55:34 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: 1992103 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=BtXGubif; 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 4XJgvQ6Wc3z1xtY for ; Thu, 3 Oct 2024 02:56:50 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B9DA1385C6D3 for ; Wed, 2 Oct 2024 16:56:48 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wr1-x435.google.com (mail-wr1-x435.google.com [IPv6:2a00:1450:4864:20::435]) by sourceware.org (Postfix) with ESMTPS id 1DA5D3858D3C for ; Wed, 2 Oct 2024 16:56:00 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 1DA5D3858D3C 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 1DA5D3858D3C Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::435 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1727888166; cv=none; b=LxXJg+Zi/+q9gNxBPCcYDDMkkkR41cEARHUD8zTAKMCYmAQ81r+9sh47DclZzXvc3Rz4w54gCKZy74a9lld5pVV3jKtQePL/ZHLUIpAL3tQn2aCJu19d4G1TO60Ov+aVEvKK8vTjDkNR5XfsrELuaNEdHSa5FHprz7PZGrulU8U= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1727888166; c=relaxed/simple; bh=NzIBd4sIBslxQmSmU4WsdiAaSTlUeA07O0BbSZV83yM=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=A1IoUfE/YSiDIa2y9uSxNaBi+8GsmtNEp6uWzpphl0yJJiNcfA/tCvi4xYG1OR43J85kK8UTvWpgmQpzM1PeD2yLrQ6CrZQ4F7DD16WGrCA0bF8B/WzdL/ewj/bsif0aTlRqz1+W9uUgdE7n+AGFc6Gvu5l4NrbHkzCcNQ0PK28= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wr1-x435.google.com with SMTP id ffacd0b85a97d-37d04b32ea6so22890f8f.1 for ; Wed, 02 Oct 2024 09:56:00 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1727888158; x=1728492958; 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=M6/uCxepgnpHKlaEOPY3zjldBiXVBGYlt8yYJAXibnU=; b=BtXGubifJBL7hpQKrLWppIKIVpmgzsZDLXFJtJuqk4LjLgu9s654y4sYL1mtYcx2CC YpVZfW8TV3tYI5xsI3u8hLeM3G4o1zpujf7a7UrNisTU3/8dq+mEeZsaI6uaj3uGHGxZ 8epHNixwt6FlrSXt0VX6hsx+fEoOLBmYnIQdV1RSn11Lq5z7xw34eDm92hb4ahqbTVwK J5L8NV5s5rMa+6+viGOgxuTcM/Z/hfWB896ckeasKG7dQlaHXXq83lJovdUJCIWlg0lL iWkb045mQ5y+GZA+OdIj0gHh85I0WYld62222O/35w3pJINPF0UZWWgNdfK5leltkFJD q3oA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1727888158; x=1728492958; 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=M6/uCxepgnpHKlaEOPY3zjldBiXVBGYlt8yYJAXibnU=; b=YiepTPbdaXZfZJUZvhsX1g6Lz2L5bo2CNkV/KuLs+g2qPTPJcfejM4t7YDI9YAUZbk zSo7S1TEnwf2RZsz2maKNyLtOC1rHAeLTE30oJthQK5jirUtSM5vEktuusheryNM4RLX R3NNJVtYp2wqdQfS7FiYrC5yabSR5u1QLAqjPBYUs8M7oSur4Lkd+q6hpvoywb/kWZP9 TE/FC/NpDG6FCGCye6XcTV6Nh4K6X5rk6SIRm8ziWDYm+ODsm7262raercQEaryhefx1 6wA+iciXq4G9CbAKWzoSekoiMHm0rHY6r08nQmBxtEZIPzuiU9x0LzwkaGBq4wiDKomq zyqg== X-Gm-Message-State: AOJu0YxKSJsSERohkJhqvJ0IjaReEpVCrvJJfluM2tXfw1hVA0sP7jNe cqMRWAP6c37DVRajn2GJRjknUo5E0WRI1a7Wv7puBXFTQ+5LIjnWIxRgLNfBapyTtA89IqFhw94 B X-Google-Smtp-Source: AGHT+IEbeKkWYqmTFKUI72V3VMbKUFDkgdenxk1u2xUvEBRYknVLQk1qBxByhvvLyRRLhItT3dZHeA== X-Received: by 2002:a5d:640f:0:b0:37c:cfbb:d356 with SMTP id ffacd0b85a97d-37cfb9d61afmr2273253f8f.28.1727888158172; Wed, 02 Oct 2024 09:55:58 -0700 (PDT) Received: from localhost.localdomain (2a01cb0018c22f0003fd7af961426ad9.ipv6.abo.wanadoo.fr. [2a01:cb00:18c2:2f00:3fd:7af9:6142:6ad9]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-37cd575de73sm14244859f8f.115.2024.10.02.09.55.57 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 02 Oct 2024 09:55:57 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: Paul-Antoine Arras Subject: [PATCH v4 3/7] OpenMP: C front-end support for dispatch + adjust_args Date: Wed, 2 Oct 2024 18:55:34 +0200 Message-ID: <20241002165538.3237107-4-parras@baylibre.com> X-Mailer: git-send-email 2.45.2 In-Reply-To: <20241002165538.3237107-1-parras@baylibre.com> References: <20241002165538.3237107-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.8 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 | 536 +++++++++++++++++++--- 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 +++ libgomp/testsuite/libgomp.c/dispatch-2.c | 84 ++++ 10 files changed, 733 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 create mode 100644 libgomp/testsuite/libgomp.c/dispatch-2.c diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc index 4dd2eecbea5..fab9b5b8b23 100644 --- a/gcc/c-family/c-attribs.cc +++ b/gcc/c-family/c-attribs.cc @@ -571,6 +571,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 620a3c1353a..5a0ed636677 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4300,8 +4300,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 ed2a7a00e9e..040370cbb6f 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 a681438cbbe..61403093382 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -1747,6 +1747,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 @@ -15090,6 +15092,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)) @@ -15098,6 +15102,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)) @@ -19365,6 +19371,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 ) */ @@ -19984,6 +20044,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; @@ -23794,6 +23862,202 @@ 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"); + + // Extract depend clauses and create taskwait + tree depend_clauses = NULL_TREE; + tree *depend_clauses_ptr = &depend_clauses; + for (tree c = OMP_DISPATCH_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) + { + *depend_clauses_ptr = c; + depend_clauses_ptr = &OMP_CLAUSE_CHAIN (c); + } + } + if (depend_clauses != NULL_TREE) + { + tree stmt = make_node (OMP_TASK); + TREE_TYPE (stmt) = void_node; + OMP_TASK_CLAUSES (stmt) = depend_clauses; + OMP_TASK_BODY (stmt) = NULL_TREE; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); + } + + // 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; + } + + // Walk the tree to find the dispatch function call and wrap it into an IFN + gcc_assert (TREE_CODE (dispatch_body) == CALL_EXPR + || TREE_CODE (dispatch_body) == MODIFY_EXPR); + tree *dispatch_call = TREE_CODE (dispatch_body) == MODIFY_EXPR + ? &TREE_OPERAND (dispatch_body, 1) + : &dispatch_body; + if (TREE_CODE (*dispatch_call) == FLOAT_EXPR + || TREE_CODE (*dispatch_call) == CONVERT_EXPR) + dispatch_call = &TREE_OPERAND (*dispatch_call, 0); + *dispatch_call = build_call_expr_internal_loc ( + loc, IFN_GOMP_DISPATCH, + TREE_TYPE (TREE_TYPE (CALL_EXPR_FN (*dispatch_call))), 1, *dispatch_call); + + 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 @@ -24774,6 +25038,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 \ @@ -25237,77 +25505,223 @@ 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; + tree need_device_ptr_list = make_node (TREE_LIST); - 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); // : + + 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) + { + need_device_ptr_list = chainon ( + need_device_ptr_list, + build_tree_list ( + NULL_TREE, + build_int_cst ( + integer_type_node, + idx))); // Store 0-based argument index, + // as in gimplify_call_expr + } + } + } + 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"); + } } - parens.require_close (parser); - c_parser_skip_to_pragma_eol (parser); + if (TREE_CHAIN (need_device_ptr_list) != NULL_TREE + && variant != error_mark_node) + { + tree variant_decl = tree_strip_nop_conversions (variant); + DECL_ATTRIBUTES (variant_decl) + = tree_cons (get_identifier ("omp declare variant variant adjust_args"), + build_tree_list (need_device_ptr_list, + NULL_TREE /*need_device_addr */), + DECL_ATTRIBUTES (variant_decl)); + } } /* Finalize #pragma omp declare simd or #pragma omp declare variant @@ -26127,7 +26541,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) \ @@ -26135,7 +26548,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) @@ -27053,6 +27470,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 ba6d96d26b2..fa4777dd35a 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -16333,6 +16333,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; +} diff --git a/libgomp/testsuite/libgomp.c/dispatch-2.c b/libgomp/testsuite/libgomp.c/dispatch-2.c new file mode 100644 index 00000000000..faa0d8a1d1c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/dispatch-2.c @@ -0,0 +1,84 @@ +// Adapted from OpenMP examples + +#include +#include +#include + +int baz (double *d_bv, const double *d_av, int n); +int bar (double *d_bv, const double *d_av, int n); + +#pragma omp declare variant(bar) match(construct={dispatch}) adjust_args(need_device_ptr: f_bv, f_av) +#pragma omp declare variant(baz) match(implementation={vendor(gnu)}) +int foo (double *f_bv, const double *f_av, int n); + +int baz (double *bv, const double *av, int n); +int bar (double *bv, const double *av, int n); + +int foo (double *bv, const double *av, int n) +{ + for (int i = 0; i < n; i++) + bv[i] = av[i] * i; + return -1; +} + +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; +} + +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; +}