From patchwork Fri Feb 16 23:35:48 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 85913 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 4CF68385773A for ; Fri, 16 Feb 2024 23:36:35 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ed1-x52c.google.com (mail-ed1-x52c.google.com [IPv6:2a00:1450:4864:20::52c]) by sourceware.org (Postfix) with ESMTPS id 6A38C3858C5E for ; Fri, 16 Feb 2024 23:35:52 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6A38C3858C5E 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 6A38C3858C5E Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::52c ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1708126559; cv=none; b=auv+5KFZB8nEQ29D5QdPksJZNXPqLLamVQxzPrLkONXjaO3ITDb17YCSkqsgMFRj8XnZ1+eKYnxkikj/hH+b316rRrSrRWTrV5z2Ab3YLIA53l7hNh1V+N3UA9/tFIPn3RC1Ii7SOX6tU/2gqqr9YThgoEKum7csmjsUYXSpcA4= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1708126559; c=relaxed/simple; bh=6QZ++K/MfM222XMIqGxiIL9dT1BSBDDU1rsAnN06f6Q=; h=DKIM-Signature:Message-ID:Date:MIME-Version:Subject:To:From; b=DVagnTFZrycK+SCapypVvYDIHgUWE8AO6gUXylhSgozJBKS3txma7GsmMMZs0CnL2WWVzCo8lZ3l8WOxvm56mvEDK/DLlplwmUSh6DQyRcMi+dmAf+ZyNwiQQgkIH0idSK01qsSrWGHCEermWnOVPOm6GDhXJX3MhU7riCHxnzA= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ed1-x52c.google.com with SMTP id 4fb4d7f45d1cf-5640fef9fa6so854563a12.0 for ; Fri, 16 Feb 2024 15:35:52 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1708126551; x=1708731351; darn=gcc.gnu.org; h=in-reply-to:from:content-language:references:to:subject:user-agent :mime-version:date:message-id:from:to:cc:subject:date:message-id :reply-to; bh=8z7DHrk8aIlkepTESMw5boFKTX7a3JtKIHUk5YECSU0=; b=KpRFDodoJLuL9HYh8NMxh5RIoOx/66w5ibDsnywn585Ubb211C1+qzu92eV3YZpRyi wd80VQAtRmQ1fnBpmxEns0at1tzrNAku76Zhn01LOrwVaFz0tdKhrftz/KrxEBe8xpm0 k3VSWDt/eP3Nua1o6RcLEuqONqZzoxNZQnRNpsWlyKFDx7weSnY+1sHI+tsviy3bR343 Afs7DrSXqUcVAYEnXJIcQ4NkF1ihkeHTR0jZEhZyae+BlErK5VMPsJ3lVPSIoNCYJCnM RHPUpKGdCyWaRCZwIb6b9qAraIshZvPcTTCWYWIid3gajWGOWMWPnpFyBKoewZEbHPAW xftg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1708126551; x=1708731351; h=in-reply-to:from:content-language:references:to:subject:user-agent :mime-version:date:message-id:x-gm-message-state:from:to:cc:subject :date:message-id:reply-to; bh=8z7DHrk8aIlkepTESMw5boFKTX7a3JtKIHUk5YECSU0=; b=qT253tokBW+JGHK+ltadXk8piphJ1/T+w6dKKJLrNN2Z3hKbcapxh6pnvjqXrvDbla iRd7TMlJuFJhm4DPwKXvdCvXRfaoto2f+8wA+7rkwNTnuP7mLltBCv4YmElrc5VHARIu KYfOP60Tzc4rD8XfBxmYpUkUGKVoGUSlOjZLWEX2WxHAwewm0srGDKhrt21iHTc+4NNU UA0RnWTOTphz5sbCUqdj1IvMZYLxg1bfT4sJ90FCPa9uJFkNCGN7D9h4zQfHxnq11FIH NFJElOFFQjyB7LoDWWmhWWzgnNXFe9C3FO4/iOBUUK37AJl+fG4JDkyBfQD55aUuWN1Y 4Kbw== X-Forwarded-Encrypted: i=1; AJvYcCU0vUQr9yz/YZ/gC4rgwmFWgCyxHHO9zjO0gGwRNKD8jVDSBKclNOtWL4VhH0wn4qk87gy1IvYKMGZwl/IdRUn8pJ2RTq0BzA== X-Gm-Message-State: AOJu0YwNyQcjXggu8X3Y7nrQCFje60PO95oSQDHpEW8CaqRIJUS4mM1Y 0AbFwYdIlZGO3OllU70jH44I68GRanfLV3smGOiwI6WVZOFmivDmzxpQslRJ88s= X-Google-Smtp-Source: AGHT+IHv454AM2Psj6Zwx2jpUR7eJBel5Q415H7F+cf19P1T93MzCLchokQ9h63iNbyrtzUTn5vE2w== X-Received: by 2002:a17:906:a446:b0:a3d:3ea1:e6e1 with SMTP id cb6-20020a170906a44600b00a3d3ea1e6e1mr4855596ejb.30.1708126550736; Fri, 16 Feb 2024 15:35:50 -0800 (PST) Received: from ?IPV6:2001:16b8:3fe3:e300:322f:d86c:a5cb:ffc8? ([2001:16b8:3fe3:e300:322f:d86c:a5cb:ffc8]) by smtp.gmail.com with ESMTPSA id qw13-20020a1709066a0d00b00a3df2b849a5sm434164ejc.155.2024.02.16.15.35.49 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 16 Feb 2024 15:35:50 -0800 (PST) Message-ID: Date: Sat, 17 Feb 2024 00:35:48 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Subject: [Patch] OpenMP/C++: Fix (first)private clause with member variables [PR110347] [was: [RFA/RFC] C++/OpenMP: Supporting (first)private for member variables [PR110347] - or VALUE_EXPR and gimplify] To: Jakub Jelinek , gcc-patches References: Content-Language: en-US From: Tobias Burnus In-Reply-To: X-Spam-Status: No, score=-10.8 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_FILL_THIS_FORM_SHORT, 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.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Hi, your suggestion almost did the trick, but caused regressions with lambda closures in target regions. Jakub Jelinek wrote: > Ah, and the reason why it doesn't work on target is that it has the > everything is mapped assumption: > if ((ctx->region_type & ORT_TARGET) != 0) > { > if (ctx->region_type & ORT_ACC) > /* For OpenACC, as remarked above, defer expansion. */ > shared = false; > else > shared = true; > > ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared); > > Perhaps shared = true; should be shared = (flags & GOVD_MAPPED) != 0; > now that we have private/firstprivate clauses on target? Hence, I now use this code, but also pass a flag to distinguish target regions (→ map) from shared usage, assuming that it is needed for the latter (otherwise, there wouldn't be that code). The issue only showed up for a compile-only testcase, which I have now turned into a run-time testcase. In order to do so, I had to fix a bogus test for is mapped (or at least I think it is bogus) - and for sure it didn't handle shared memory. I also modified it such that it iterates over devices. Changes to the dump: the 'device' clause had to be added (3x) and for the long line: 'this' and 'iptr' swapped the order and 'map(from:mapped)' became 'firstprivate(mapped)' due to my changes. I appended a patch which only shows the test-case differences as "git diff" contains all lines as I move it to libgomp/. Comments, remarks, suggestions? Tobias OpenMP/C++: Fix (first)private clause with member variables [PR110347] OpenMP permits '(first)private' for C++ member variables, which GCC handles by tagging those by DECL_OMP_PRIVATIZED_MEMBER, adding a temporary VAR_DECL and DECL_VALUE_EXPR pointing to the 'this->member_var' in the C++ front end. The idea is that in omp-low.cc, the DECL_VALUE_EXPR is used before the region (for 'firstprivate'; ignored for 'private') while in the region, the DECL itself is used. In gimplify, the value expansion is suppressed and deferred if the lang_hooks.decls.omp_disregard_value_expr (decl, shared) returns true - which is never the case if 'shared' is true. In OpenMP 4.5, only 'map' and 'use_device_ptr' was permitted for the 'target' directive. And when OpenMP 5.0's 'private'/'firstprivate' clauses was added, the update that 'shared' is only true for 'map' was missed. However, just enabling it for all '!shared' will cause issues with Lambda closures ("__closure->this->...") for which also a DECL_VALUE_EXPR exists but that is not related to DECL_OMP_PRIVATIZED_MEMBER. Solution: Update the lang hook to take a Boolean argument, indicating whether it is called for a target region or not. 2024-02-16 Tobias Burnus Jakub Jelinek PR c++/110347 gcc/cp/ChangeLog: * cp-gimplify.cc (cxx_omp_disregard_value_expr): Add new Boolean argument and use it. * cp-tree.h (cxx_omp_disregard_value_expr): Update prototype. gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_omp_disregard_value_expr): Add unused Boolean argument. * trans.h (gfc_omp_disregard_value_expr): Update prototype. gcc/ChangeLog: * gimplify.cc (omp_notice_variable): Update call to lang_hooks.decls.omp_disregard_value_expr. (omp_notice_variable): Likewise; fix 'shared' arg for (first)private in target regions. * hooks.cc (hook_bool_tree_bool_bool_false): New. * hooks.h (hook_bool_tree_bool_bool_false): New. * langhooks-def.h (LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR): Use it. * langhooks.h (struct lang_hooks_for_decls): Add second Boolean argument. * omp-low.cc (omp_member_access_dummy_var): Update lang_hooks.decls.omp_disregard_value_expr call. libgomp/ChangeLog: * testsuite/libgomp.c++/target-lambda-3.C: Moved from gcc/testsuite/g++.dg/gomp/ and fixed is-mapped handling. * testsuite/libgomp.c++/firstprivate-c++-1.C: New test. * testsuite/libgomp.c++/firstprivate-c++-2.C: New test. * testsuite/libgomp.c++/private-c++-1.C: New test. * testsuite/libgomp.c++/private-c++-2.C: New test. * testsuite/libgomp.c++/use_device_ptr-c++-1.C: New test. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-lambda-1.C: Moved to become a run-time test under testsuite/libgomp.c++. Co-authored-by: Jakub Jelinek gcc/cp/cp-gimplify.cc | 7 +- gcc/cp/cp-tree.h | 2 +- gcc/fortran/trans-openmp.cc | 2 +- gcc/fortran/trans.h | 2 +- gcc/gimplify.cc | 12 +- gcc/hooks.cc | 6 + gcc/hooks.h | 1 + gcc/langhooks-def.h | 2 +- gcc/langhooks.h | 5 +- gcc/omp-low.cc | 2 +- gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 94 ------- libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C | 305 +++++++++++++++++++++ libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C | 125 +++++++++ libgomp/testsuite/libgomp.c++/private-c++-1.C | 247 +++++++++++++++++ libgomp/testsuite/libgomp.c++/private-c++-2.C | 117 ++++++++ libgomp/testsuite/libgomp.c++/target-lambda-3.C | 104 +++++++ .../testsuite/libgomp.c++/use_device_ptr-c++-1.C | 125 +++++++++ 17 files changed, 1048 insertions(+), 110 deletions(-) diff --git a/gcc/cp/cp-gimplify.cc b/gcc/cp/cp-gimplify.cc index 30e94797f9f..dcc46d86619 100644 --- a/gcc/cp/cp-gimplify.cc +++ b/gcc/cp/cp-gimplify.cc @@ -2754,10 +2754,11 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */) /* Return true if DECL's DECL_VALUE_EXPR (if any) should be disregarded in OpenMP construct, because it is going to be remapped during OpenMP lowering. SHARED is true if DECL - is going to be shared, false if it is going to be privatized. */ + is going to be shared, false if it is going to be privatized. TARGET is + true if this for an OpenMP target/OpenACC compute region. */ bool -cxx_omp_disregard_value_expr (tree decl, bool shared) +cxx_omp_disregard_value_expr (tree decl, bool shared, bool target) { if (shared) return false; @@ -2767,7 +2768,7 @@ cxx_omp_disregard_value_expr (tree decl, bool shared) && DECL_LANG_SPECIFIC (decl) && DECL_OMP_PRIVATIZED_MEMBER (decl)) return true; - if (VAR_P (decl) && DECL_CONTEXT (decl) && is_capture_proxy (decl)) + if (!target && VAR_P (decl) && DECL_CONTEXT (decl) && is_capture_proxy (decl)) return true; return false; } diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 334c11396c2..2dc200cd43e 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -8464,7 +8464,7 @@ extern tree cxx_omp_clause_assign_op (tree, tree, tree); extern tree cxx_omp_clause_dtor (tree, tree); extern void cxx_omp_finish_clause (tree, gimple_seq *, bool); extern bool cxx_omp_privatize_by_reference (const_tree); -extern bool cxx_omp_disregard_value_expr (tree, bool); +extern bool cxx_omp_disregard_value_expr (tree, bool, bool); extern void cp_fold_function (tree); extern tree cp_fold_maybe_rvalue (tree, bool); extern tree cp_fold_rvalue (tree); diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index a2bf15665b3..74e213ab09e 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -1791,7 +1791,7 @@ gfc_omp_scalar_target_p (tree decl) is going to be shared, false if it is going to be privatized. */ bool -gfc_omp_disregard_value_expr (tree decl, bool shared) +gfc_omp_disregard_value_expr (tree decl, bool shared, bool /* target */) { if (GFC_DECL_COMMON_OR_EQUIV (decl) && DECL_HAS_VALUE_EXPR_P (decl)) diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h index 2e10ce1a9b3..d8e640ade27 100644 --- a/gcc/fortran/trans.h +++ b/gcc/fortran/trans.h @@ -840,7 +840,7 @@ void gfc_omp_finish_clause (tree, gimple_seq *, bool); bool gfc_omp_allocatable_p (tree); bool gfc_omp_scalar_p (tree, bool); bool gfc_omp_scalar_target_p (tree); -bool gfc_omp_disregard_value_expr (tree, bool); +bool gfc_omp_disregard_value_expr (tree, bool, bool); bool gfc_omp_private_debug_clause (tree, bool); bool gfc_omp_private_outer_ref (tree); struct gimplify_omp_ctx; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 7f79b3cc7e6..dc524dc12b0 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -7990,7 +7990,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) the whole block. For C++ and Fortran, it can also be true under certain other conditions, if DECL_HAS_VALUE_EXPR. */ if (RECORD_OR_UNION_TYPE_P (type)) - is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false); + is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false, true); if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0 && is_global_var (decl) @@ -8092,7 +8092,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) } if (ctx->region_type == ORT_NONE) - return lang_hooks.decls.omp_disregard_value_expr (decl, false); + return lang_hooks.decls.omp_disregard_value_expr (decl, false, false); if (is_global_var (decl)) { @@ -8148,9 +8148,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) /* For OpenACC, as remarked above, defer expansion. */ shared = false; else - shared = true; + shared = (flags & GOVD_MAP) != 0; - ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared); + ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared, true); if (n == NULL) { unsigned nflags = flags; @@ -8305,7 +8305,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) omp_add_variable (ctx, decl, flags); shared = (flags & GOVD_SHARED) != 0; - ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared); + ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared, false); goto do_outer; } @@ -8350,7 +8350,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) shared = false; else shared = ((flags | n->value) & GOVD_SHARED) != 0; - ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared); + ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared, false); /* If nothing changed, there's nothing left to do. */ if ((n->value & flags) == flags) diff --git a/gcc/hooks.cc b/gcc/hooks.cc index 28769074222..32eccfd5f08 100644 --- a/gcc/hooks.cc +++ b/gcc/hooks.cc @@ -343,6 +343,12 @@ hook_bool_tree_bool_false (tree, bool) return false; } +bool +hook_bool_tree_bool_bool_false (tree, bool, bool) +{ + return false; +} + bool hook_bool_rtx_insn_true (rtx_insn *) { diff --git a/gcc/hooks.h b/gcc/hooks.h index 924748420e6..02dc63dc3c1 100644 --- a/gcc/hooks.h +++ b/gcc/hooks.h @@ -72,6 +72,7 @@ extern bool hook_bool_rtx_mode_int_int_intp_bool_false (rtx, machine_mode, extern bool hook_bool_tree_tree_false (tree, tree); extern bool hook_bool_tree_tree_true (tree, tree); extern bool hook_bool_tree_bool_false (tree, bool); +extern bool hook_bool_tree_bool_bool_false (tree, bool, bool); extern bool hook_bool_wint_wint_uint_bool_true (const widest_int &, const widest_int &, unsigned int, bool); diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h index f5c67b6823c..67c100a0af3 100644 --- a/gcc/langhooks-def.h +++ b/gcc/langhooks-def.h @@ -263,7 +263,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree); #define LANG_HOOKS_OMP_PREDETERMINED_SHARING lhd_omp_predetermined_sharing #define LANG_HOOKS_OMP_PREDETERMINED_MAPPING lhd_omp_predetermined_mapping #define LANG_HOOKS_OMP_REPORT_DECL lhd_pass_through_t -#define LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR hook_bool_tree_bool_false +#define LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR hook_bool_tree_bool_bool_false #define LANG_HOOKS_OMP_PRIVATE_DEBUG_CLAUSE hook_bool_tree_bool_false #define LANG_HOOKS_OMP_PRIVATE_OUTER_REF hook_bool_tree_false #define LANG_HOOKS_OMP_CLAUSE_DEFAULT_CTOR hook_tree_tree_tree_tree_null diff --git a/gcc/langhooks.h b/gcc/langhooks.h index 5a4dfb6ef62..68bd91f3c62 100644 --- a/gcc/langhooks.h +++ b/gcc/langhooks.h @@ -280,8 +280,9 @@ struct lang_hooks_for_decls /* Return true if DECL's DECL_VALUE_EXPR (if any) should be disregarded in OpenMP construct, because it is going to be remapped during OpenMP lowering. SHARED is true if DECL - is going to be shared, false if it is going to be privatized. */ - bool (*omp_disregard_value_expr) (tree, bool); + is going to be shared, false if it is going to be privatized. TARGET + is true when this if for an OpenMP target/OPenACC compute contruct. */ + bool (*omp_disregard_value_expr) (tree, bool, bool); /* Return true if DECL that is shared iff SHARED is true should be put into OMP_CLAUSE_PRIVATE_DEBUG. */ diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 4d003f42098..7b4631029c7 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -275,7 +275,7 @@ omp_member_access_dummy_var (tree decl) || !DECL_ARTIFICIAL (decl) || !DECL_IGNORED_P (decl) || !DECL_HAS_VALUE_EXPR_P (decl) - || !lang_hooks.decls.omp_disregard_value_expr (decl, false)) + || !lang_hooks.decls.omp_disregard_value_expr (decl, false, false)) return NULL_TREE; tree v = DECL_VALUE_EXPR (decl); diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C deleted file mode 100644 index 5ce8ceadb19..00000000000 --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C +++ /dev/null @@ -1,94 +0,0 @@ -// We use 'auto' without a function return type, so specify dialect here -// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } -#include -#include - -template -void -omp_target_loop (int begin, int end, L loop) -{ - #pragma omp target teams distribute parallel for - for (int i = begin; i < end; i++) - loop (i); -} - -struct S -{ - int a, len; - int *ptr; - - auto merge_data_func (int *iptr, int &b) - { - auto fn = [=](void) -> bool - { - bool mapped; - #pragma omp target map(from:mapped) - { - mapped = (ptr != NULL && iptr != NULL); - if (mapped) - { - for (int i = 0; i < len; i++) - ptr[i] += a + b + iptr[i]; - } - } - return mapped; - }; - return fn; - } -}; - -int x = 1; - -int main (void) -{ - const int N = 10; - int *data1 = new int[N]; - int *data2 = new int[N]; - memset (data1, 0xab, sizeof (int) * N); - memset (data1, 0xcd, sizeof (int) * N); - - int val = 1; - int &valref = val; - #pragma omp target enter data map(alloc: data1[:N], data2[:N]) - - omp_target_loop (0, N, [=](int i) { data1[i] = val; }); - omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }); - - #pragma omp target update from(data1[:N], data2[:N]) - - for (int i = 0; i < N; i++) - { - if (data1[i] != 1) abort (); - if (data2[i] != 2) abort (); - } - - #pragma omp target exit data map(delete: data1[:N], data2[:N]) - - int b = 8; - S s = { 4, N, data1 }; - auto f = s.merge_data_func (data2, b); - - if (f ()) abort (); - - #pragma omp target enter data map(to: data1[:N]) - if (f ()) abort (); - - #pragma omp target enter data map(to: data2[:N]) - if (!f ()) abort (); - - #pragma omp target exit data map(from: data1[:N], data2[:N]) - - for (int i = 0; i < N; i++) - { - if (data1[i] != 0xf) abort (); - if (data2[i] != 2) abort (); - } - - return 0; -} - -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ - -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */ - -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C b/libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C new file mode 100644 index 00000000000..ae5d4fbe1bf --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/firstprivate-c++-1.C @@ -0,0 +1,305 @@ +/* PR c++/110347 */ + +#include +#include +#include + +struct S { + int A, B[10], *C; + void f (int dev); + void g (int dev); +}; + +template +struct St { + T A, B[10], *C; + void ft (int dev); + void gt (int dev); +}; + + +void +S::f (int dev) +{ + A = 5; + C = (int *) malloc (sizeof (int) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \ + firstprivate(c_saved) device(dev) + { + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) + { + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + +void +S::g (int dev) +{ + A = 5; + C = (int *) malloc (sizeof (int) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \ + device(dev) + { +#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */ + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); +#endif + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) + { + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + + +template +void +St::ft (int dev) +{ + A = 5; + C = (T *) malloc (sizeof (T) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \ + firstprivate(c_saved) device(dev) + { + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) + { + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + +template +void +St::gt (int dev) +{ + A = 5; + C = (T *) malloc (sizeof (T) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target firstprivate(A) firstprivate(B) firstprivate(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \ + device(dev) + { +#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */ + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); +#endif + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) firstprivate(A) firstprivate(B) firstprivate(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) + { + if (A != 5) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + +int +main () +{ + struct S s; + struct St st; + for (int dev = 0; dev <= omp_get_num_devices(); dev++) + { + s.f (dev); + st.ft (dev); + s.g (dev); + st.gt (dev); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C b/libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C new file mode 100644 index 00000000000..a4f2514b591 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/firstprivate-c++-2.C @@ -0,0 +1,125 @@ +/* PR c++/110347 */ + +#include + +struct t { + int A; + void f (int dev); +}; + +void +t::f (int dev) +{ + int B = 49; + + A = 7; + #pragma omp parallel firstprivate(A) if(0) shared(B) default(none) + { + if (A != 7) { __builtin_printf("ERROR 1b: %d (!= 7) inside omp parallel\n", A); __builtin_abort (); } + A = 5; + B = A; + } + if (A != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", A); __builtin_abort (); } + if (B != 5) { __builtin_printf("ERROR 1a: %d\n", B); __builtin_abort (); } + A = 8; B = 49; + #pragma omp parallel firstprivate(A)if(0) shared(B) default(none) + { + if (A != 8) { __builtin_printf("ERROR 1b: %d (!= 8) inside omp parallel\n", A); __builtin_abort (); } + A = 6; + B = A; + } + if (A != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", A); __builtin_abort (); } + if (B != 6) { __builtin_printf("ERROR 2a: %d\n", B); __builtin_abort (); } + A = 8; B = 49; + + #pragma omp target firstprivate(A) map(from:B) device(dev) + { + if (A != 8) { __builtin_printf("ERROR 2b: %d (!= 8) inside omp target\n", A); __builtin_abort (); } + A = 7; + B = A; + } + if (A != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", A); __builtin_abort (); } + if (B != 7) { __builtin_printf("ERROR 3a: %d\n", B); __builtin_abort (); } + A = 9; B = 49; + #pragma omp target firstprivate(A) map(from:B) device(dev) + { + if (A != 9) { __builtin_printf("ERROR 3b: %d (!= 9) inside omp target\n", A); __builtin_abort (); } + A = 8; + B = A; + } + if (A != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", A); __builtin_abort (); } + if (B != 8) { __builtin_printf("ERROR 4a: %d\n", B); __builtin_abort (); } +} + + +template +struct tt { + T C; + void g (int dev); +}; + +template +void +tt::g (int dev) +{ + T D = 49; + C = 7; + #pragma omp parallel firstprivate(C) if(0) shared(D) default(none) + { + if (C != 7) { __builtin_printf("ERROR 1b: %d (!= 7) inside omp parallel\n", C);__builtin_abort (); } + C = 5; + D = C; + } + if (C != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", C);__builtin_abort (); } + if (D != 5) { __builtin_printf("ERROR 1a: %d\n", D);__builtin_abort (); } + C = 8; D = 49; + #pragma omp parallel firstprivate(C)if(0) shared(D) default(none) + { + if (C != 8) { __builtin_printf("ERROR 1b: %d (!= 8) inside omp parallel\n", C);__builtin_abort (); } + C = 6; + D = C; + } + if (C != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", C);__builtin_abort (); } + if (D != 6) { __builtin_printf("ERROR 2a: %d\n", D);__builtin_abort (); } + C = 8; D = 49; + #pragma omp target firstprivate(C) map(from:D) defaultmap(none) device(dev) + { + if (C != 8) { __builtin_printf("ERROR 2b: %d (!= 8) inside omp target\n", C);__builtin_abort (); } + C = 7; + D = C; + } + if (C != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", C);__builtin_abort (); } + if (D != 7) { __builtin_printf("ERROR 3a: %d\n", D);__builtin_abort (); } + C = 9; D = 49; + #pragma omp target firstprivate(C) map(from:D) defaultmap(none) device(dev) + { + if (C != 9) { __builtin_printf("ERROR 3b: %d (!= 9) inside omp target\n", C);__builtin_abort (); } + C = 8; + D = C; + } + if (C != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", C); __builtin_abort (); } + if (D != 8) { __builtin_printf("ERROR 4a: %d\n", D); } +} + +void +foo () +{ + struct t x; + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + x.f (dev); +} + +void +bar () +{ + struct tt y; + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + y.g (dev); +} + +int +main () +{ + foo (); + bar (); +} diff --git a/libgomp/testsuite/libgomp.c++/private-c++-1.C b/libgomp/testsuite/libgomp.c++/private-c++-1.C new file mode 100644 index 00000000000..19ee726a222 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/private-c++-1.C @@ -0,0 +1,247 @@ +/* PR c++/110347 */ + +#include +#include +#include + +struct S { + int A, B[10], *C; + void f (int dev); + void g (int dev); +}; + +template +struct St { + T A, B[10], *C; + void ft (int dev); + void gt (int dev); +}; + + +void +S::f (int dev) +{ + A = 5; + C = (int *) malloc (sizeof (int) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target private(A) private(B) private(C) device(dev) + { + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) private(A) private(B) private(C) + { + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + +void +S::g (int dev) +{ + A = 5; + C = (int *) malloc (sizeof (int) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target private(A) private(B) private(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \ + device(dev) + { +#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */ + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); +#endif + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) private(A) private(B) private(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) + { + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + + +template +void +St::ft (int dev) +{ + A = 5; + C = (T *) malloc (sizeof (T) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target private(A) private(B) private(C) device(dev) + { + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) private(A) private(B) private(C) + { + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + +template +void +St::gt (int dev) +{ + A = 5; + C = (T *) malloc (sizeof (T) * 10); + uintptr_t c_saved = (uintptr_t) C; + for (int i = 0; i < 10; i++) + B[i] = C[i] = i+5; + + #pragma omp target private(A) private(B) private(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) \ + device(dev) + { +#if 0 /* FIXME: The following is disabled because of PR middle-end/113436. */ + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); +#endif + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + #pragma omp parallel if (0) private(A) private(B) private(C) \ + allocate(allocator(omp_low_lat_mem_alloc), align(128): A, B, C) + { + if (((uintptr_t) &A) % 128 != 0) + abort (); + if (((uintptr_t) &B) % 128 != 0) + abort (); + if (((uintptr_t) &C) % 128 != 0) + abort (); + A = 99; + for (int i = 0; i < 10; i++) + B[i] = -i-23; + C = &A; + } + + if (A != 5) + abort (); + if (c_saved != (uintptr_t) C) + abort (); + for (int i = 0; i < 10; i++) + if (B[i] != i + 5 || C[i] != i+5) + abort (); + + free (C); +} + +int +main () +{ + struct S s; + struct St st; + for (int dev = 0; dev <= omp_get_num_devices(); dev++) + { + s.f (dev); + st.ft (dev); + s.g (dev); + st.gt (dev); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/private-c++-2.C b/libgomp/testsuite/libgomp.c++/private-c++-2.C new file mode 100644 index 00000000000..aa472cb62ee --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/private-c++-2.C @@ -0,0 +1,117 @@ +/* PR c++/110347 */ + +#include + +struct t { + int A; + void f (int dev); +}; + +void +t::f (int dev) +{ + int B = 49; + + A = 7; + #pragma omp parallel private(A) if(0) shared(B) default(none) + { + A = 5; + B = A; + } + if (A != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", A); __builtin_abort (); } + if (B != 5) { __builtin_printf("ERROR 1a: %d\n", B); __builtin_abort (); } + A = 8; B = 49; + #pragma omp parallel private(A)if(0) shared(B) default(none) + { + A = 6; + B = A; + } + if (A != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", A); __builtin_abort (); } + if (B != 6) { __builtin_printf("ERROR 2a: %d\n", B); __builtin_abort (); } + A = 8; B = 49; + + #pragma omp target private(A) map(from:B) device(dev) + { + A = 7; + B = A; + } + if (A != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", A); __builtin_abort (); } + if (B != 7) { __builtin_printf("ERROR 3a: %d\n", B); __builtin_abort (); } + A = 9; B = 49; + #pragma omp target private(A) map(from:B) device(dev) + { + A = 8; + B = A; + } + if (A != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", A); __builtin_abort (); } + if (B != 8) { __builtin_printf("ERROR 4a: %d\n", B); __builtin_abort (); } +} + + +template +struct tt { + T C; + void g (int dev); +}; + +template +void +tt::g (int dev) +{ + T D = 49; + C = 7; + #pragma omp parallel private(C) if(0) shared(D) default(none) + { + C = 5; + D = C; + } + if (C != 7) { __builtin_printf("ERROR 1: %d (!= 7) omp parallel\n", C);__builtin_abort (); } + if (D != 5) { __builtin_printf("ERROR 1a: %d\n", D);__builtin_abort (); } + C = 8; D = 49; + #pragma omp parallel private(C)if(0) shared(D) default(none) + { + C = 6; + D = C; + } + if (C != 8) { __builtin_printf("ERROR 2: %d (!= 8) omp parallel\n", C);__builtin_abort (); } + if (D != 6) { __builtin_printf("ERROR 2a: %d\n", D);__builtin_abort (); } + C = 8; D = 49; + #pragma omp target private(C) map(from:D) defaultmap(none) device(dev) + { + C = 7; + D = C; + } + if (C != 8) { __builtin_printf("ERROR 3: %d (!= 8) omp target\n", C);__builtin_abort (); } + if (D != 7) { __builtin_printf("ERROR 3a: %d\n", D);__builtin_abort (); } + C = 9; D = 49; + #pragma omp target private(C) map(from:D) defaultmap(none) device(dev) + { + C = 8; + D = C; + } + if (C != 9) { __builtin_printf("ERROR 4: %d (!= 9) omp target\n", C); __builtin_abort (); } + if (D != 8) { __builtin_printf("ERROR 4a: %d\n", D); } +} + +void +foo () +{ + struct t x; + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + x.f (dev); +} + +void +bar () +{ + struct tt y; + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + y.g (dev); +} + +int +main () +{ + foo (); + bar (); +} diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-3.C b/libgomp/testsuite/libgomp.c++/target-lambda-3.C new file mode 100644 index 00000000000..6be8426bd3e --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-lambda-3.C @@ -0,0 +1,104 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include +#include +#include + +template +void +omp_target_loop (int begin, int end, L loop, int dev) +{ + #pragma omp target teams distribute parallel for device(dev) + for (int i = begin; i < end; i++) + loop (i); +} + +struct S +{ + int a, len; + int *ptr; + + auto merge_data_func (int *iptr, int &b, int dev) + { + auto fn = [=](void) -> bool + { + bool mapped = (omp_target_is_present (iptr, dev) + && omp_target_is_present (ptr, dev)); + #pragma omp target device(dev) + { + if (mapped) + { + for (int i = 0; i < len; i++) + ptr[i] += a + b + iptr[i]; + } + } + return mapped; + }; + return fn; + } +}; + +int x = 1; + +void run (int dev) +{ + const int N = 10; + int *data1 = new int[N]; + int *data2 = new int[N]; + memset (data1, 0xab, sizeof (int) * N); + memset (data2, 0xcd, sizeof (int) * N); + + bool shared_mem = (omp_target_is_present (data1, dev) + && omp_target_is_present (data2, dev)); + int val = 1; + int &valref = val; + #pragma omp target enter data map(alloc: data1[:N], data2[:N]) device(dev) + + omp_target_loop (0, N, [=](int i) { data1[i] = val; }, dev); + omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }, dev); + + #pragma omp target update from(data1[:N], data2[:N]) device(dev) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 1) abort (); + if (data2[i] != 2) abort (); + } + + #pragma omp target exit data map(delete: data1[:N], data2[:N]) device(dev) + + int b = 8; + S s = { 4, N, data1 }; + auto f = s.merge_data_func (data2, b, dev); + if (f () ^ shared_mem) abort (); + + #pragma omp target enter data map(to: data1[:N]) device(dev) + if (f () ^ shared_mem) abort (); + + #pragma omp target enter data map(to: data2[:N]) device(dev) + if (!f ()) abort (); + + #pragma omp target exit data map(from: data1[:N], data2[:N]) device(dev) + + for (int i = 0; i < N; i++) + { + if ((!shared_mem && data1[i] != 0xf) + || (shared_mem && data1[i] != 0x2b)) + abort (); + if (data2[i] != 2) abort (); + } + delete [] data1; + delete [] data2; +} + +int main () +{ + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + run (dev); +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) firstprivate\(mapped\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(_[0-9]+\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) device\(dev.[0-9_]+\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\)} "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C b/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C new file mode 100644 index 00000000000..0bb6ce6434b --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/use_device_ptr-c++-1.C @@ -0,0 +1,125 @@ +/* PR c++/110347 */ + +#include + +#define N 30 + +struct t { + int *A; + void f (int dev); +}; + +void +t::f (int dev) +{ + int *ptr; + int B[N]; + for (int i = 0; i < N; i++) + B[i] = 1 + i; + ptr = A = (int *) omp_target_alloc (sizeof (int) * N, dev); + omp_target_memcpy (A, B, sizeof (int) * N, 0, 0, dev, omp_initial_device); + + #pragma omp target is_device_ptr (A) device(dev) + { + for (int i = 0; i < N; i++) + if (A[i] != 1 + i) + __builtin_abort (); + for (int i = 0; i < N; i++) + A[i] = (-2-i)*10; + A = (int *) 0x12345; + } + if (ptr != A) + __builtin_abort (); + + #pragma omp target is_device_ptr (A) device(dev) + { + for (int i = 0; i < N; i++) + if (A[i] != (-2-i)*10) + __builtin_abort (); + for (int i = 0; i < N; i++) + A[i] = (3+i)*11; + A = (int *) 0x12345; + } + if (ptr != A) + __builtin_abort (); + + int *C = (int *) __builtin_malloc (sizeof(int)*N); + omp_target_memcpy (C, A, sizeof (int) * N, 0, 0, omp_initial_device, dev); + for (int i = 0; i < N; i++) + if (C[i] != (3+i)*11) + __builtin_abort (); + __builtin_free (C); + omp_target_free (A, dev); +} + +template +struct tt { + T *D; + void g (int dev); +}; + +template +void +tt::g (int dev) +{ + T *ptr; + T E[N]; + for (int i = 0; i < N; i++) + E[i] = 1 + i; + ptr = D = (T *) omp_target_alloc (sizeof (T) * N, dev); + omp_target_memcpy (D, E, sizeof (T) * N, 0, 0, dev, omp_initial_device); + + #pragma omp target is_device_ptr (D) device(dev) + { + for (int i = 0; i < N; i++) + if (D[i] != 1 + i) + __builtin_abort (); + for (int i = 0; i < N; i++) + D[i] = (-2-i)*10; + D = (T *) 0x12345; + } + if (ptr != D) + __builtin_abort (); + + #pragma omp target is_device_ptr (D) device(dev) + { + for (int i = 0; i < N; i++) + if (D[i] != (-2-i)*10) + __builtin_abort (); + for (int i = 0; i < N; i++) + D[i] = (3+i)*11; + D = (T *) 0x12345; + } + if (ptr != D) + __builtin_abort (); + + T *F = (T *) __builtin_malloc (sizeof(T)*N); + omp_target_memcpy (F, D, sizeof (T) * N, 0, 0, omp_initial_device, dev); + for (int i = 0; i < N; i++) + if (F[i] != (3+i)*11) + __builtin_abort (); + __builtin_free (F); + omp_target_free (D, dev);} + +void +foo () +{ + struct t x; + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + x.f (dev); +} + +void +bar () +{ + struct tt y; + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + y.g (dev); +} + +int +main () +{ + foo (); + bar (); +}