From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 13A4B3858C52; Fri, 14 Jul 2023 16:34:07 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 13A4B3858C52 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="6.01,206,1684828800"; d="scan'208";a="11870706" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 14 Jul 2023 08:34:03 -0800 IronPort-SDR: ZDpazwYaqV0KiXtdO1Dz3m1aImLiUUJ4imxhvcgsWnQQHXTv4BmpQh6bqoOyZvg+B7ptN7NQXZ YPVzOkI3JewZdQsbXfXAmpOaHGoT0n1haUIngFVlfEXO3Yx1pLNFYTKEdVLWOM9tC6BjZlLsIP 5jA9+T/yhLnl755ol88m+TVG24KVy5wLX0G1EdZtGh25tI0mtvGDJ/FPJ4/YFEg+OMOjEdtZoJ fML6CLJCLDmtR9p4gpwdgffFH/z2fwUx2DiVpNNARBXw+Ug4TwDk9S8I1Tl6IddnN6bpzGNiUv RXc= From: Julian Brown To: CC: , , Subject: [PATCH] [og13] OpenMP: Dimension ordering for array-shaping operator for C and C++ Date: Fri, 14 Jul 2023 16:33:39 +0000 Message-ID: <20230714163340.3464603-1-julian@codesourcery.com> X-Mailer: git-send-email 2.25.1 MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain 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.8 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,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 List-Id: This patch fixes a bug in non-contiguous 'target update' operations using the new array-shaping operator for C and C++, processing the dimensions of the array the wrong way round during the OpenMP lowering pass. Fortran was also incorrectly using the wrong ordering but the second reversal in omp-low.cc made it produce the correct result. The C and C++ bug only affected array shapes where the dimension sizes are different ([X][Y]) - several existing tests used the same value for both/all dimensions ([X][X]), which masked the problem. Only the array dimensions (extents) are affected, not e.g. the indices, lengths or strides for array sections. This patch reverses the order used in both omp-low.cc and the Fortran front-end, so the order should now be correct for all supported base languages. Tested with offloading to AMD GCN. I will apply (to og13) shortly. 2023-07-14 Julian Brown gcc/fortran/ * trans-openmp.cc (gfc_trans_omp_arrayshape_type): Reverse dimension ordering for created array type. gcc/ * omp-low.cc (lower_omp_target): Reverse iteration over array dimensions. libgomp/ * testsuite/libgomp.c-c++-common/array-shaping-14.c: New test. --- gcc/fortran/trans-openmp.cc | 2 +- gcc/omp-low.cc | 6 ++-- .../libgomp.c-c++-common/array-shaping-14.c | 34 +++++++++++++++++++ 3 files changed, 38 insertions(+), 4 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 6cb5340687e..6b9a0430eba 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4271,7 +4271,7 @@ gfc_trans_omp_arrayshape_type (tree type, vec *dims) { gcc_assert (dims->length () > 0); - for (int i = dims->length () - 1; i >= 0; i--) + for (unsigned i = 0; i < dims->length (); i++) { tree dim = fold_convert (sizetype, (*dims)[i]); /* We need the index of the last element, not the array size. */ diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index c7706a5921f..ab2e4145ab2 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -14290,7 +14290,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) dims++; } - int tdim = tdims.length () - 1; + unsigned tdim = 0; vec *vdim; vec *vindex; @@ -14365,7 +14365,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) nc = nc2; } - if (tdim >= 0) + if (tdim < tdims.length ()) { /* We have an array shape -- use that to find the total size of the data on the target to look up @@ -14403,7 +14403,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) "for array"); dim = index = len = stride = error_mark_node; } - tdim--; + tdim++; c = nc; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c new file mode 100644 index 00000000000..4ca6f794f93 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-14.c @@ -0,0 +1,34 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +#include +#include +#include + +typedef struct { + int *ptr; +} S; + +int main(void) +{ + S q; + q.ptr = (int *) calloc (9 * 11, sizeof (int)); + +#pragma omp target enter data map(to: q.ptr, q.ptr[0:9*11]) + +#pragma omp target + for (int i = 0; i < 9*11; i++) + q.ptr[i] = i; + +#pragma omp target update from(([9][11]) q.ptr[3:3:2][1:4:3]) + + for (int j = 0; j < 9; j++) + for (int i = 0; i < 11; i++) + if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0 + && i >= 1 && i <= 10 && ((i - 1) % 3) == 0) + assert (q.ptr[j * 11 + i] == j * 11 + i); + else + assert (q.ptr[j * 11 + i] == 0); + +#pragma omp target exit data map(release: q.ptr, q.ptr[0:9*11]) + return 0; +} -- 2.25.1