From patchwork Fri May 12 12:02:47 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 93144 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:b0ea:0:b0:3b6:4342:cba0 with SMTP id b10csp5046027vqo; Fri, 12 May 2023 05:04:10 -0700 (PDT) X-Google-Smtp-Source: ACHHUZ6MuSzfIhaNooTC8b9Bh6TDSl7g75JDPrKNMK0tp8CVi+n3jCWASqEHCdTp+1cavS2EpNmb X-Received: by 2002:a50:ef05:0:b0:50d:682d:d431 with SMTP id m5-20020a50ef05000000b0050d682dd431mr19617740eds.16.1683893050194; Fri, 12 May 2023 05:04:10 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1683893050; cv=none; d=google.com; s=arc-20160816; b=Am2wTHeMqJlI6D8NSMoc50jMGFcZ/4wtSmvXFubvfDcqC0zAWDHbLgjCPfsIYhFsLg jg8Dt0nLDRuoBEoxiTFFf0KmLAfkDAQqFsF98KnMl9Dp0S1iqf5H1WwIE1FLPjfmQMpo MSfjKueAhRChn4KFsNjA83eoqNN+9jAmOpGOu0CSj743KRjI/DN0UwIhW6h1wwoTjlop MnSETlwkcCZudKxPhhMEgY+hxGR9wOs1llPCr7ukXtaPLs9OVd+Ev1ejbMQ8mFMMKlSx GY6XmdW/YQc3PpJz8e0GLx2avYd6Lb+zlYg72tbVCsMjKPWaNxIdvAbHolzxy2GY0Vm8 BwDg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:message-id:date:subject:cc:to:from:ironport-sdr :dmarc-filter:delivered-to; bh=0q6Brvu5naLNxmjf2dIa9zn9Y8JXN1jTUfqGkMyNXmw=; b=hDMS3dYn0/P7vyHI8XP3kZxYum475umvzJNxRrb+7P6cMtwMVaoOqr/yxYHkcjir3Q qWPakJf2czZpDiZ2TmZ/HC4vaLUSP0jG8p6kWuZ3P3HCjpGUTS54IDoqBmqW2PHZOW7N DeKAk4jJv8TDIA1bHr+xi6TfM3sq5SshUAMLCXw7uXGK7JlG1s06sfl1DAoKVhajCnI0 OgbzddHvKp9wD+6TWVZvZHI6TxX6prcLtgyMJVmQq9LWaGMJO9cvr3EVyccBdQoPxm7K ScOKNFZHBeYwDU6goS5+yrgDPHNNmciqiNAcrksOsnPXGzbTCp87a4FU2MEBSZ8rN21J xcMg== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id j18-20020aa7ca52000000b00506883845c3si6290709edt.623.2023.05.12.05.04.09 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 12 May 2023 05:04:10 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 258AD3855582 for ; Fri, 12 May 2023 12:04:02 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id CD9DE3858C78 for ; Fri, 12 May 2023 12:03:27 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org CD9DE3858C78 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.99,269,1677571200"; d="scan'208";a="5041272" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 12 May 2023 04:03:22 -0800 IronPort-SDR: GYiTpTfd4UeQc6HlRjYEsSn660H+Sq/65+ME60VabQOi89Lh2CkbaURaMwSLT6wnYKO/oJm4/e LhuA5+YfNwD6Mc9DvXGxSeSdfM20q2pximbcLGpDhQzmYx4whw8FP/xZjzF9qIGK/McwCUe5QT 2VuQ+AM7PYNLxUzXvTnPmbP6fCAVILoH6oCapUiL0aMLB04XzkJU39j6kwJ8EYK5haM+KPKgMW E5rCxZQSkCcpj1AkjTbqeJOac53AYwYSL3KvMjMjyY5KX2fu2vnuMJo8Ovb3WOYbmGXZMhc1Kb 374= From: Julian Brown To: CC: , , Kwok Cheung Yeung Subject: [PATCH] OpenMP: Constructors and destructors for "declare target" static aggregates Date: Fri, 12 May 2023 12:02:47 +0000 Message-ID: <20230512120247.3213280-1-julian@codesourcery.com> X-Mailer: git-send-email 2.25.1 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE 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.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1765689839064958656?= X-GMAIL-MSGID: =?utf-8?q?1765689839064958656?= This patch adds support for running constructors and destructors for static (file-scope) aggregates for C++ objects which are marked with "declare target" directives on OpenMP offload targets. At present, space is allocated on the target for such aggregates, but nothing ever constructs them properly, so they end up zero-initialised. The approach taken is to generate a set of constructors to run on the target: this currently works for AMD GCN, but fails on NVPTX due to lack of constructor/destructor support there so far on mainline. (See the new test static-aggr-constructor-destructor-3.C for a reason why running constructors on the target is preferable to e.g. constructing on the host and then copying the resulting object to the target.) This patch was previously posted for the og12 branch here: https://gcc.gnu.org/pipermail/gcc-patches/2023-March/614710.html https://gcc.gnu.org/pipermail/gcc-patches/2023-April/615013.html https://gcc.gnu.org/pipermail/gcc-patches/2023-April/615144.html though needed a fair amount of rework for mainline due to Nathan's (earlier!) patch: https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596402.html Tested with offloading to AMD GCN and bootstrapped. OK for mainline? Thanks, Julian 2023-05-12 Julian Brown gcc/cp/ * decl2.cc (tree-inline.h): Include. (static_init_fini_fns): Bump to four entries. Update comment. (start_objects, start_partial_init_fini_fn): Add 'omp_target' parameter. Support "declare target" decls. Update forward declaration. (emit_partial_init_fini_fn): Add 'host_fn' parameter. Return tree for the created function. Support "declare target". (OMP_SSDF_IDENTIFIER): New macro. (partition_vars_for_init_fini): Support partitioning "declare target" variables also. (generate_ctor_or_dtor_function): Add 'omp_target' parameter. Support "declare target" decls. (c_parse_final_cleanups): Support constructors/destructors on OpenMP offload targets. gcc/ * omp-builtins.def (BUILT_IN_OMP_IS_INITIAL_DEVICE): New builtin. * tree.cc (get_file_function_name): Support names for on-target constructor/destructor functions. libgomp/ * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test. --- gcc/cp/decl2.cc | 243 +++++++++++++++--- gcc/omp-builtins.def | 2 + gcc/tree.cc | 6 +- .../static-aggr-constructor-destructor-1.C | 28 ++ .../static-aggr-constructor-destructor-2.C | 31 +++ .../static-aggr-constructor-destructor-3.C | 36 +++ 6 files changed, 305 insertions(+), 41 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C create mode 100644 libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C create mode 100644 libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc index b510cdac554..ceec681fbeb 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -50,20 +50,22 @@ along with GCC; see the file COPYING3. If not see #include "asan.h" #include "optabs-query.h" #include "omp-general.h" +#include "tree-inline.h" /* Id for dumping the raw trees. */ int raw_dump_id; extern cpp_reader *parse_in; -static tree start_objects (bool, unsigned, bool); +static tree start_objects (bool, unsigned, bool, bool); static tree finish_objects (bool, unsigned, tree, bool = true); -static tree start_partial_init_fini_fn (bool, unsigned, unsigned); +static tree start_partial_init_fini_fn (bool, unsigned, unsigned, bool); static void finish_partial_init_fini_fn (tree); -static void emit_partial_init_fini_fn (bool, unsigned, tree, - unsigned, location_t); +static tree emit_partial_init_fini_fn (bool, unsigned, tree, + unsigned, location_t, tree); static void one_static_initialization_or_destruction (bool, tree, tree); -static void generate_ctor_or_dtor_function (bool, unsigned, tree, location_t); +static void generate_ctor_or_dtor_function (bool, unsigned, tree, location_t, + bool); static tree prune_vars_needing_no_initialization (tree *); static void write_out_vars (tree); static void import_export_class (tree); @@ -165,9 +167,10 @@ struct priority_map_traits typedef hash_map priority_map_t; -/* A pair of such hash tables, indexed by initp -- one for fini and - one for init. The fini table is only ever used when !cxa_atexit. */ -static GTY(()) priority_map_t *static_init_fini_fns[2]; +/* Two pairs of such hash tables, for the host and an OpenMP offload device. + Each pair has one priority map for fini and one for init. The fini tables + are only ever used when !cxa_atexit. */ +static GTY(()) priority_map_t *static_init_fini_fns[4]; /* Nonzero if we're done parsing and into end-of-file activities. */ @@ -3867,7 +3870,8 @@ generate_tls_wrapper (tree fn) /* Start a global constructor or destructor function. */ static tree -start_objects (bool initp, unsigned priority, bool has_body) +start_objects (bool initp, unsigned priority, bool has_body, + bool omp_target = false) { bool default_init = initp && priority == DEFAULT_INIT_PRIORITY; bool is_module_init = default_init && module_global_init_needed (); @@ -3881,7 +3885,15 @@ start_objects (bool initp, unsigned priority, bool has_body) /* We use `I' to indicate initialization and `D' to indicate destruction. */ - unsigned len = sprintf (type, "sub_%c", initp ? 'I' : 'D'); + unsigned len; + if (omp_target) + /* Use "off_" signifying "offload" here. The name must be distinct + from the non-offload case. The format of the name is scanned in + tree.cc/get_file_function_name, so stick to the same length for + both name variants. */ + len = sprintf (type, "off_%c", initp ? 'I' : 'D'); + else + len = sprintf (type, "sub_%c", initp ? 'I' : 'D'); if (priority != DEFAULT_INIT_PRIORITY) { char joiner = '_'; @@ -3896,6 +3908,17 @@ start_objects (bool initp, unsigned priority, bool has_body) tree fntype = build_function_type (void_type_node, void_list_node); tree fndecl = build_lang_decl (FUNCTION_DECL, name, fntype); + + if (omp_target) + { + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), NULL_TREE, + DECL_ATTRIBUTES (fndecl)); + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target nohost"), NULL_TREE, + DECL_ATTRIBUTES (fndecl)); + } + DECL_CONTEXT (fndecl) = FROB_CONTEXT (global_namespace); if (is_module_init) { @@ -3980,34 +4003,63 @@ finish_objects (bool initp, unsigned priority, tree body, bool startp) /* The name of the function we create to handle initializations and destructions for objects with static storage duration. */ #define SSDF_IDENTIFIER "__static_initialization_and_destruction" +#define OMP_SSDF_IDENTIFIER "__omp_target_static_init_and_destruction" /* Begins the generation of the function that will handle all initialization or destruction of objects with static storage duration at PRIORITY. - It is assumed that this function will only be called once. */ + It is assumed that this function will be called once for the host, and once + for an OpenMP offload target. */ static tree -start_partial_init_fini_fn (bool initp, unsigned priority, unsigned count) +start_partial_init_fini_fn (bool initp, unsigned priority, unsigned count, + bool omp_target) { - char id[sizeof (SSDF_IDENTIFIER) + 1 /* '\0' */ + 32]; + tree name; - /* Create the identifier for this function. It will be of the form - SSDF_IDENTIFIER_. */ - sprintf (id, "%s_%u", SSDF_IDENTIFIER, count); + if (omp_target) + { + char id[sizeof (OMP_SSDF_IDENTIFIER) + 1 /* \0 */ + 32]; + + /* Create the identifier for this function. It will be of the form + OMP_SSDF_IDENTIFIER_. */ + sprintf (id, "%s_%u", OMP_SSDF_IDENTIFIER, count); + name = get_identifier (id); + } + else + { + char id[sizeof (SSDF_IDENTIFIER) + 1 /* '\0' */ + 32]; + /* Create the identifier for this function. It will be of the form + SSDF_IDENTIFIER_. */ + sprintf (id, "%s_%u", SSDF_IDENTIFIER, count); + name = get_identifier (id); + } tree type = build_function_type (void_type_node, void_list_node); /* Create the FUNCTION_DECL itself. */ - tree fn = build_lang_decl (FUNCTION_DECL, get_identifier (id), type); + tree fn = build_lang_decl (FUNCTION_DECL, name, type); TREE_PUBLIC (fn) = 0; DECL_ARTIFICIAL (fn) = 1; + if (omp_target) + { + DECL_ATTRIBUTES (fn) + = tree_cons (get_identifier ("omp declare target"), NULL_TREE, + DECL_ATTRIBUTES (fn)); + DECL_ATTRIBUTES (fn) + = tree_cons (get_identifier ("omp declare target nohost"), NULL_TREE, + DECL_ATTRIBUTES (fn)); + } + + int idx = initp + 2 * omp_target; + /* Put this function in the list of functions to be called from the static constructors and destructors. */ - if (!static_init_fini_fns[initp]) - static_init_fini_fns[initp] = priority_map_t::create_ggc (); - auto &slot = static_init_fini_fns[initp]->get_or_insert (priority); + if (!static_init_fini_fns[idx]) + static_init_fini_fns[idx] = priority_map_t::create_ggc (); + auto &slot = static_init_fini_fns[idx]->get_or_insert (priority); slot = tree_cons (fn, NULL_TREE, slot); /* Put the function in the global scope. */ @@ -4203,22 +4255,76 @@ one_static_initialization_or_destruction (bool initp, tree decl, tree init) a TREE_LIST of VAR_DECL with static storage duration. Whether initialization or destruction is performed is specified by INITP. */ -static void +static tree emit_partial_init_fini_fn (bool initp, unsigned priority, tree vars, - unsigned counter, location_t locus) + unsigned counter, location_t locus, tree host_fn) { input_location = locus; - tree body = start_partial_init_fini_fn (initp, priority, counter); + bool omp_target = (host_fn != NULL_TREE); + tree body = start_partial_init_fini_fn (initp, priority, counter, omp_target); + tree fndecl = current_function_decl; + + tree nonhost_if_stmt = NULL_TREE; + if (omp_target) + { + nonhost_if_stmt = begin_if_stmt (); + /* We add an "omp declare target nohost" attribute, but (for + now) we still get a copy of the constructor/destructor on + the host. Make sure it does nothing unless we're on the + target device. */ + tree fn + = builtin_decl_explicit (BUILT_IN_OMP_IS_INITIAL_DEVICE); + tree initial_dev = build_call_expr (fn, 0); + tree target_dev_p + = cp_build_binary_op (input_location, NE_EXPR, initial_dev, + build_int_cst (NULL_TREE, 1), + tf_warning_or_error); + finish_if_stmt_cond (target_dev_p, nonhost_if_stmt); + } for (tree node = vars; node; node = TREE_CHAIN (node)) - /* Do one initialization or destruction. */ - one_static_initialization_or_destruction (initp, TREE_VALUE (node), - TREE_PURPOSE (node)); + { + tree decl = TREE_VALUE (node); + tree init = TREE_PURPOSE (node); + /* We will emit 'init' twice, and it is modified in-place during + gimplification. Make a copy here. */ + if (omp_target) + { + /* We've already emitted INIT in the host version of the ctor/dtor + function. We need to deep-copy it (including new versions of + local variables introduced, etc.) for use in the target + ctor/dtor function. */ + copy_body_data id; + hash_map decl_map; + memset (&id, 0, sizeof (id)); + id.src_fn = host_fn; + id.dst_fn = current_function_decl; + id.src_cfun = DECL_STRUCT_FUNCTION (id.src_fn); + id.decl_map = &decl_map; + id.copy_decl = copy_decl_no_change; + id.transform_call_graph_edges = CB_CGE_DUPLICATE; + id.transform_new_cfg = true; + id.transform_return_to_modify = false; + id.eh_lp_nr = 0; + walk_tree (&init, copy_tree_body_r, &id, NULL); + } + /* Do one initialization or destruction. */ + one_static_initialization_or_destruction (initp, decl, init); + } + + if (omp_target) + { + /* Finish up nonhost if-stmt body. */ + finish_then_clause (nonhost_if_stmt); + finish_if_stmt (nonhost_if_stmt); + } /* Finish up the static storage duration function for this round. */ input_location = locus; finish_partial_init_fini_fn (body); + + return fndecl; } /* VARS is a list of variables with static storage duration which may @@ -4281,7 +4387,7 @@ prune_vars_needing_no_initialization (tree *vars) This reverses the variable ordering. */ void -partition_vars_for_init_fini (tree var_list, priority_map_t *(&parts)[2]) +partition_vars_for_init_fini (tree var_list, priority_map_t *(&parts)[4]) { for (auto node = var_list; node; node = TREE_CHAIN (node)) { @@ -4307,6 +4413,30 @@ partition_vars_for_init_fini (tree var_list, priority_map_t *(&parts)[2]) auto &slot = parts[false]->get_or_insert (priority); slot = tree_cons (NULL_TREE, decl, slot); } + + if (flag_openmp + && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + { + priority_map_t **omp_parts = parts + 2; + + if (init || (flag_use_cxa_atexit && has_cleanup)) + { + // Add to initialization list. + if (!omp_parts[true]) + omp_parts[true] = priority_map_t::create_ggc (); + auto &slot = omp_parts[true]->get_or_insert (priority); + slot = tree_cons (init, decl, slot); + } + + if (!flag_use_cxa_atexit && has_cleanup) + { + // Add to finalization list. + if (!omp_parts[false]) + omp_parts[false] = priority_map_t::create_ggc (); + auto &slot = omp_parts[false]->get_or_insert (priority); + slot = tree_cons (NULL_TREE, decl, slot); + } + } } } @@ -4334,10 +4464,10 @@ write_out_vars (tree vars) static void generate_ctor_or_dtor_function (bool initp, unsigned priority, - tree fns, location_t locus) + tree fns, location_t locus, bool omp_target) { input_location = locus; - tree body = start_objects (initp, priority, bool (fns)); + tree body = start_objects (initp, priority, bool (fns), omp_target); if (fns) { @@ -4979,7 +5109,7 @@ c_parse_final_cleanups (void) auto_vec consteval_vtables; int retries = 0; - unsigned ssdf_count = 0; + unsigned ssdf_count = 0, omp_ssdf_count = 0; for (bool reconsider = true; reconsider; retries++) { reconsider = false; @@ -5042,8 +5172,9 @@ c_parse_final_cleanups (void) write_out_vars (vars); function_depth++; // Disable GC - priority_map_t *parts[2] = {nullptr, nullptr}; + priority_map_t *parts[4] = {nullptr, nullptr, nullptr, nullptr}; partition_vars_for_init_fini (vars, parts); + tree host_init_fini[2] = { NULL_TREE, NULL_TREE }; for (unsigned initp = 2; initp--;) if (parts[initp]) @@ -5054,10 +5185,32 @@ c_parse_final_cleanups (void) // Partitioning kept the vars in reverse order. // We only want that for dtors. list = nreverse (list); - emit_partial_init_fini_fn (initp, iter.first, list, - ssdf_count++, - locus_at_end_of_parsing); + host_init_fini[initp] + = emit_partial_init_fini_fn (initp, iter.first, list, + ssdf_count++, + locus_at_end_of_parsing, + NULL_TREE); } + + if (flag_openmp) + { + priority_map_t **omp_parts = parts + 2; + for (unsigned initp = 2; initp--;) + if (omp_parts[initp]) + for (auto iter : *omp_parts[initp]) + { + auto list = iter.second; + if (initp) + // Partitioning kept the vars in reverse order. + // We only want that for dtors. + list = nreverse (list); + emit_partial_init_fini_fn (initp, iter.first, list, + omp_ssdf_count++, + locus_at_end_of_parsing, + host_init_fini[initp]); + } + } + function_depth--; // Re-enable GC /* All those initializations and finalizations might cause @@ -5223,7 +5376,11 @@ c_parse_final_cleanups (void) if (static_init_fini_fns[true]) for (auto iter : *static_init_fini_fns[true]) iter.second = nreverse (iter.second); - + + if (flag_openmp && static_init_fini_fns[2 + true]) + for (auto iter : *static_init_fini_fns[2 + true]) + iter.second = nreverse (iter.second); + /* Then, do the Objective-C stuff. This is where all the Objective-C module stuff gets generated (symtab, class/protocol/selector lists etc). This must be done after C++ @@ -5238,7 +5395,7 @@ c_parse_final_cleanups (void) { input_location = locus_at_end_of_parsing; tree body = start_partial_init_fini_fn (true, DEFAULT_INIT_PRIORITY, - ssdf_count++); + ssdf_count++, false); /* For Objective-C++, we may need to initialize metadata found in this module. This must be done _before_ any other static initializations. */ @@ -5257,18 +5414,26 @@ c_parse_final_cleanups (void) static_init_fini_fns[true] = priority_map_t::create_ggc (); if (static_init_fini_fns[true]->get_or_insert (DEFAULT_INIT_PRIORITY)) has_module_inits = true; + + if (flag_openmp) + { + if (!static_init_fini_fns[2 + true]) + static_init_fini_fns[2 + true] = priority_map_t::create_ggc (); + static_init_fini_fns[2 + true]->get_or_insert (DEFAULT_INIT_PRIORITY); + } } /* Generate initialization and destruction functions for all priorities for which they are required. They have C-language linkage. */ push_lang_context (lang_name_c); - for (unsigned initp = 2; initp--;) + for (unsigned initp = 4; initp--;) if (static_init_fini_fns[initp]) { for (auto iter : *static_init_fini_fns[initp]) - generate_ctor_or_dtor_function (initp, iter.first, iter.second, - locus_at_end_of_parsing); + generate_ctor_or_dtor_function (initp & 1, iter.first, iter.second, + locus_at_end_of_parsing, + (initp & 2) != 0); static_init_fini_fns[initp] = nullptr; } pop_lang_context (); diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index e0f03263db0..a80158f5c14 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -68,6 +68,8 @@ DEF_GOACC_BUILTIN_ONLY (BUILT_IN_GOACC_SINGLE_COPY_START, "GOACC_single_copy_sta DEF_GOACC_BUILTIN_ONLY (BUILT_IN_GOACC_SINGLE_COPY_END, "GOACC_single_copy_end", BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_OMP_IS_INITIAL_DEVICE, "omp_is_initial_device", + BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_THREAD_NUM, "omp_get_thread_num", BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_THREADS, "omp_get_num_threads", diff --git a/gcc/tree.cc b/gcc/tree.cc index 7e6de288886..097806e9875 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -8806,9 +8806,11 @@ get_file_function_name (const char *type) will be local to this file and the name is only necessary for debugging purposes. We also assign sub_I and sub_D sufixes to constructors called from - the global static constructors. These are always local. */ + the global static constructors. These are always local. + OpenMP "declare target" offloaded constructors/destructors use "off_I" and + "off_D" for the same purpose. */ else if (((type[0] == 'I' || type[0] == 'D') && targetm.have_ctors_dtors) - || (startswith (type, "sub_") + || ((startswith (type, "sub_") || startswith (type, "off_")) && (type[4] == 'I' || type[4] == 'D'))) { const char *file = main_input_filename; diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C new file mode 100644 index 00000000000..91d8469a150 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C @@ -0,0 +1,28 @@ +// { dg-do run } + +#include + +#pragma omp declare target + +struct str { + str(int x) : _x(x) { } + int add(str o) { return _x + o._x; } + int _x; +} v1(5); + +#pragma omp end declare target + +int main() +{ + int res = -1; + str v2(2); + +#pragma omp target map(from:res) + { + res = v1.add(v2); + } + + assert (res == 7); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C new file mode 100644 index 00000000000..1bf3ee8e31c --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C @@ -0,0 +1,31 @@ +// { dg-do run } + +#include + +#pragma omp declare target + +template +struct str { + str(T x) : _x(x) { } + T add(str o) { return _x + o._x; } + T _x; +}; + +str v1(5); + +#pragma omp end declare target + +int main() +{ + long res = -1; + str v2(2); + +#pragma omp target map(from:res) + { + res = v1.add(v2); + } + + assert (res == 7); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C new file mode 100644 index 00000000000..8d4aff21cd7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C @@ -0,0 +1,36 @@ +// { dg-do run } + +#include + +#pragma omp declare target + +struct item { + item(item *p, int v) : prev(p), val(v) { } + int get() { return prev ? prev->get() * val : val; } + item *prev; + int val; +}; + +/* This case demonstrates why constructing on the host and then copying to + the target would be less desirable. With on-target construction, "prev" + for each 'item' will be a device pointer, not a host pointer. */ +item hubert1(nullptr, 3); +item hubert2(&hubert1, 5); +item hubert3(&hubert2, 7); +item hubert4(&hubert3, 11); + +#pragma omp end declare target + +int main() +{ + int res = -1; + +#pragma omp target map(from:res) + { + res = hubert4.get (); + } + + assert (res == 1155); + + return 0; +}