From patchwork Sun May 26 16:46:19 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1105595 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-501683-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 45BmGl0gM0z9s4V for ; Mon, 27 May 2019 02:46:49 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type; q=dns; s= default; b=rOqmKvNryXiITDhVlZR6eZBUfGT+hLvbZ5WxJOpBGKCgLZ3cv5HcA K6ZSzGROiGWfYppTgmNe9defsk5I1AcQTS1WAMFSyP+PZVdlINq6fHMyzJOsuGxa UA1DT+nWqMnhaw7gDMY/v79At+grmNRrnp64FD1dqqJ/qiRVe6mCS4= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type; s= default; bh=WMSCc4ZklX34SewpYmX2GRHfTbo=; b=UKucPxZi+ekYGX7RNeJ2 wLzMXea+7zgLhETCBGQxM7Uxzhzrlowu0CyJhQUqTHkxKcA6EzoiVct07e9ETpfE gZVrlHpCxvBvXxuHimawbvDEL+TEmmV7qReQWMzndqPNBnACH9ncuKwjRircSTfB vY+CGQjj6yEV/WOYtz7PGc8= Received: (qmail 82022 invoked by alias); 26 May 2019 16:46:40 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 82014 invoked by uid 89); 26 May 2019 16:46:40 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-15.8 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_PASS autolearn=ham version=3.3.1 spammy=recorded, doubt, tree_static, TREE_STATIC X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sun, 26 May 2019 16:46:38 +0000 Received: from svr-orw-mbx-06.mgc.mentorg.com ([147.34.90.206]) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1hUwI3-0000WL-6U from Thomas_Schwinge@mentor.com ; Sun, 26 May 2019 09:46:31 -0700 Received: from SVR-ORW-MBX-07.mgc.mentorg.com (147.34.90.207) by SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Sun, 26 May 2019 09:46:28 -0700 Received: from tftp-cs (147.34.91.1) by SVR-ORW-MBX-07.mgc.mentorg.com (147.34.90.207) with Microsoft SMTP Server id 15.0.1320.4 via Frontend Transport; Sun, 26 May 2019 09:46:28 -0700 Received: by tftp-cs (Postfix, from userid 49978) id 2AA87C23E9; Sun, 26 May 2019 09:46:28 -0700 (PDT) From: Thomas Schwinge To: , Jakub Jelinek Subject: Make possible 'scan-tree-dump' of 'lower_omp_target' mapping kinds User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Sun, 26 May 2019 18:46:19 +0200 Message-ID: <8736l1qi3o.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! To establish some suitable testsuite coverage for a task that I'm working on, I need to do 'scan-tree-dump' of 'lower_omp_target' mapping kinds. Is the attached OK? Any suggestions about whether/how to restrict the (effective?) targets this gets run for, because no doubt there are target-specific bits at least in the alignment chosen. The attached test case passes for x86_64-pc-linux-gnu with '--target_board=unix\{,-m32,-mx32\}'. (I didn't verify the mappings generated, but just documented the status quo.) If approving this patch, please respond with "Reviewed-by: NAME " so that your effort will be recorded in the commit log, see . Grüße Thomas From ba19ef0d5aebf2604e6625d1cca81fc477801966 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Sun, 26 May 2019 18:33:26 +0200 Subject: [PATCH] Make possible 'scan-tree-dump' of 'lower_omp_target' mapping kinds gcc/ * omp-low.c (lower_omp_target): Dump mapping kinds. gcc/testsuite/ * c-c++-common/gomp/lower_omp_target-mappings-1.c: New file. --- gcc/omp-low.c | 103 +++++++++++++++--- .../gomp/lower_omp_target-mappings-1.c | 49 +++++++++ 2 files changed, 136 insertions(+), 16 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c diff --git a/gcc/omp-low.c b/gcc/omp-low.c index faab5d384280..79468289d88f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -9218,15 +9218,23 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) static void lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { + pretty_printer pp; tree clauses; tree child_fn, t, c; gomp_target *stmt = as_a (gsi_stmt (*gsi_p)); gbind *tgt_bind, *bind, *dep_bind = NULL; gimple_seq tgt_body, olist, ilist, fplist, new_body; location_t loc = gimple_location (stmt); + const char *loc_str = NULL; bool offloaded, data_region; unsigned int map_cnt = 0; + if (dump_file && (dump_flags & TDF_DETAILS)) + { + dump_location (&pp, loc); + loc_str = pp_formatted_text (&pp); + } + offloaded = is_gimple_omp_offloaded (stmt); switch (gimple_omp_target_kind (stmt)) { @@ -9687,7 +9695,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (TREE_CODE (s) != INTEGER_CST) TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0; - unsigned HOST_WIDE_INT tkind, tkind_zero; + unsigned HOST_WIDE_INT tkind, tkind_align; + unsigned HOST_WIDE_INT tkind_zero, tkind_zero_align; switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_MAP: @@ -9743,25 +9752,43 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) < (HOST_WIDE_INT_C (1U) << talign_shift)); gcc_checking_assert (tkind_zero < (HOST_WIDE_INT_C (1U) << talign_shift)); - talign = ceil_log2 (talign); - tkind |= talign << talign_shift; - tkind_zero |= talign << talign_shift; - gcc_checking_assert (tkind + { + unsigned int talign2 = ceil_log2 (talign); + tkind_align = tkind | (talign2 << talign_shift); + tkind_zero_align = tkind_zero | (talign2 << talign_shift); + } + gcc_checking_assert (tkind_align <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); - gcc_checking_assert (tkind_zero + gcc_checking_assert (tkind_zero_align <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); - if (tkind == tkind_zero) - x = build_int_cstu (tkind_type, tkind); + if (tkind_align == tkind_zero_align) + x = build_int_cstu (tkind_type, tkind_align); else { TREE_STATIC (TREE_VEC_ELT (t, 2)) = 0; x = build3 (COND_EXPR, tkind_type, fold_build2 (EQ_EXPR, boolean_type_node, unshare_expr (s), size_zero_node), - build_int_cstu (tkind_type, tkind_zero), - build_int_cstu (tkind_type, tkind)); + build_int_cstu (tkind_type, tkind_zero_align), + build_int_cstu (tkind_type, tkind_align)); } CONSTRUCTOR_APPEND_ELT (vkind, purpose, x); + + if (dump_file && (dump_flags & TDF_DETAILS)) + { + fprintf (dump_file, "Mapping %s%s [%u] '", + loc_str, + IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)), + (unsigned int) tree_to_uhwi (purpose)); + print_generic_expr (dump_file, ovar, dump_flags); + fprintf (dump_file, "': size = "); + print_generic_expr (dump_file, s, dump_flags); + fprintf (dump_file, ", kind = %u/%u, align = %u\n", + (unsigned int) tkind, + (unsigned int) tkind_zero, + talign); + } + if (nc && nc != c) c = nc; break; @@ -9826,12 +9853,29 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_checking_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift)); - talign = ceil_log2 (talign); - tkind |= talign << talign_shift; - gcc_checking_assert (tkind + { + unsigned int talign2 = ceil_log2 (talign); + tkind_align = tkind | (talign2 << talign_shift); + } + gcc_checking_assert (tkind_align <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); CONSTRUCTOR_APPEND_ELT (vkind, purpose, - build_int_cstu (tkind_type, tkind)); + build_int_cstu (tkind_type, tkind_align)); + + if (dump_file && (dump_flags & TDF_DETAILS)) + { + fprintf (dump_file, "Mapping %s%s [%u] '", + loc_str, + IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)), + (unsigned int) tree_to_uhwi (purpose)); + print_generic_expr (dump_file, ovar, dump_flags); + fprintf (dump_file, "': size = "); + print_generic_expr (dump_file, s, dump_flags); + fprintf (dump_file, ", kind = %u, align = %u\n", + (unsigned int) tkind, + talign); + } + break; case OMP_CLAUSE_USE_DEVICE_PTR: @@ -9862,14 +9906,41 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); gcc_checking_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift)); - gcc_checking_assert (tkind + tkind_align = tkind; + gcc_checking_assert (tkind_align <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); CONSTRUCTOR_APPEND_ELT (vkind, purpose, - build_int_cstu (tkind_type, tkind)); + build_int_cstu (tkind_type, tkind_align)); + + if (dump_file && (dump_flags & TDF_DETAILS)) + { + fprintf (dump_file, "Mapping %s%s [%u] '", + loc_str, + IDENTIFIER_POINTER (DECL_NAME (ctx->cb.dst_fn)), + (unsigned int) tree_to_uhwi (purpose)); + print_generic_expr (dump_file, ovar, dump_flags); + fprintf (dump_file, "': size = "); + print_generic_expr (dump_file, s, dump_flags); + fprintf (dump_file, ", kind = %u, align = %u\n", + (unsigned int) tkind, + talign); + } + break; } gcc_assert (map_idx == map_cnt); + if (flag_checking) + for (unsigned int i = 0; i < map_cnt; ++i) + { + tree t_index = (*vsize)[i].index; + unsigned HOST_WIDE_INT ii = tree_to_uhwi (t_index); + gcc_assert (ii == i); + + t_index = (*vsize)[i].index; + ii = tree_to_uhwi (t_index); + gcc_assert (ii == i); + } DECL_INITIAL (TREE_VEC_ELT (t, 1)) = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize); diff --git a/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c b/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c new file mode 100644 index 000000000000..2e8a3064078f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/lower_omp_target-mappings-1.c @@ -0,0 +1,49 @@ +// { dg-additional-options "-fdump-tree-omplower-details" } + +int main(int argc, char *argv[]) +{ + char vla[argc + 3]; +#pragma omp target map(from:vla[0:argc]) + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[0\] '\*vla.1\[0\]': size = .*, kind = 2/15, align = 1$} 1 omplower } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[1\] '.*': size = 0, kind = 13, align = 4$} 1 omplower { target { ! { c++ && lp64 } } } } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[1\] '.*': size = 0, kind = 13, align = 8$} 1 omplower { target { c++ && lp64 } } } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[2\] 'argc': size = 0, kind = 13, align = 4$} 1 omplower } } + // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:6:9\] main\._omp_fn\.0 \[3\]} omplower } } + { + for (int i = 0; i < argc; ++i) + vla[i] = i + sizeof vla; + } + +#pragma omp target data use_device_ptr(argv) + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:17:9\] main \[0\] 'argv': size = 0, kind = 14, align = 4$} 1 omplower } } + // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:17:9\] main \[1\]} omplower } } + { +#pragma omp target map(from:argc) + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[0\] 'argc': size = 4, kind = 2/2, align = 4$} 1 omplower } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[1\] '\*vla\.1': size = .*, kind = 3/3, align = 1$} 1 omplower } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[2\] 'MEM\[\(char \*\)argv\]': size = 0, kind = 15/15, align = 1$} 1 omplower } } + // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:21:9\] main\._omp_fn\.1 \[3\]} omplower } } + { + argc = (argv != (void *) vla); + } + } + +#pragma omp target data map(from:argc) + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:31:9\] main \[0\] 'argc': size = 4, kind = 2/2, align = 4$} 1 omplower } } + // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:31:9\] main \[1\]} omplower } } + { +#pragma omp target is_device_ptr(argv) + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[0\] 'argv': size = 0, kind = 13, align = 0$} 1 omplower { target c } } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[0\] 'argv': size = 0, kind = 13, align = 1$} 1 omplower { target c++ } } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[1\] '\*vla\.1': size = .*, kind = 3/3, align = 1$} 1 omplower } } + // { dg-final { scan-tree-dump-times {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[2\] 'argc': size = 0, kind = 13, align = 4$} 1 omplower } } + // { dg-final { scan-tree-dump-not {(?n)^Mapping \[.*lower_omp_target-mappings-1.c:35:9\] main\._omp_fn\.2 \[3\]} omplower } } + { + argc = (argv != (void *) &vla[1]); + } + } + + return 0; +} + +// { dg-final { scan-tree-dump-times {(?n)^Mapping} 11 omplower } } -- 2.17.1