[0/5] openmp: Support array-of-pointers (multi-segment) noncontiguous array sections and array-shaping casts
| Message ID | 20260805202343.2868178-1-parras@baylibre.com |
|---|---|
| Headers |
Return-Path: <gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org> X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from vm01.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 15F584BB3BB8 for <patchwork@sourceware.org>; Wed, 5 Aug 2026 20:26:17 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 15F584BB3BB8 Authentication-Results: sourceware.org; dkim=pass (2048-bit key, secure) header.d=baylibre.com header.i=@baylibre.com header.a=rsa-sha256 header.s=google header.b=LsTmw6Y+ X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wr1-x42e.google.com (mail-wr1-x42e.google.com [IPv6:2a00:1450:4864:20::42e]) by sourceware.org (Postfix) with ESMTPS id 75C164BAE7DC for <gcc-patches@gcc.gnu.org>; Wed, 5 Aug 2026 20:25:14 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 75C164BAE7DC 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 75C164BAE7DC Authentication-Results: sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::42e ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961514; cv=none; b=pI1Pi1tU2ksXasiHVT3b1QimljgETpEgl+8utH3YazIpcYZFwYSMwUIPKqO6yffb65UX3Wift0cu7sm6xhFnrSFmcuk71tRvfJMP8WmFlFxlynCVGdrCRytXs5hHNjzAdzMg7XUGuAGuU1W5XIPlIyE1GS0oJyme3FUc4CCZW4c= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961514; c=relaxed/simple; bh=INYjsbOTi9O3AeqccV2MlaQ4VhHiIS3oIHc5OpJjrKI=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=OnGjQ8vz/OfKjzCXmiIqQBBet09mhXrs0Jmcl3SQ5LxMlChDKCQo30Ca0ydl1aDo7/KmmFcULL5gb9dIqYCvqIcpdDtDx5/RBW+gvTWV8cr1DXCYRpewkydjVn46353KJ279ByyD02p+NFzRRwwzfovP0h88Jhw7tijsTAk85lk= ARC-Authentication-Results: i=1; sourceware.org; dkim=pass (2048-bit key, secure) header.d=baylibre.com header.i=@baylibre.com header.a=rsa-sha256 header.s=google header.b=LsTmw6Y+ DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 75C164BAE7DC Received: by mail-wr1-x42e.google.com with SMTP id ffacd0b85a97d-47fecbbafdaso454881f8f.0 for <gcc-patches@gcc.gnu.org>; Wed, 05 Aug 2026 13:25:14 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1785961513; x=1786566313; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:message-id:date:subject:cc :to:from:from:to:cc:subject:date:message-id:reply-to:content-type; bh=2DiBuXLImENYuRRegjnd5Oeu2BvI7SctjlM/6cDpXZs=; b=LsTmw6Y+2LQF8LJdWKlUgADg57gclII+PlvicSk+DFKtxCXNk+u+qsll1kIiHvcj+z 5t4MU7PNiMT+tyUTb6NvNElcs6GiJwLMLhlZ5xvpN/dwXf/xTw/bOkghKVeWTf8LWMi2 Uw7JIVy2WpG6wN1LFUgzVTUWieghBCceZAfMhzWCziASNHghAGSvVb0gSsdADbVORdOJ B7CyJrlMk1vQdXOarKUhWn9rJZced1At/5Ahns5KGBdcqK1vEitOFwdLDTndPsUaeLPc PYxy23Rd80QiczL7tPcx+vmanyqO8QiXZA/OTwR7K2tfug+UikjNuAdv1aCCBP+E2rj1 idxA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1785961513; x=1786566313; h=content-transfer-encoding:mime-version:message-id:date:subject:cc :to:from:x-gm-gg:x-gm-message-state:from:to:cc:subject:date :message-id:reply-to:content-type; bh=2DiBuXLImENYuRRegjnd5Oeu2BvI7SctjlM/6cDpXZs=; b=oZhC4UsGINmmJPrKo26PIT7/41zq81TPtsFPauRqIkOCwbLlP/X/yexmS/rEJgVz1c MQc5BsvR8D4aaUHIJt4QTIPoD/KzQz/9qmZ67g8bnAh/NFe7k6YAy+ZUwlE4uMh8I6T4 RAAMn2ldMrc9d4XT1rcQjTAREeFvTqeZTCrVvCkBjISllP3IRZxgG1hA3lRixW5YKW5p D5qYX9PioYYGU1u6vttnLTX2vwuCLHbbXqAh6BBZMm6I0ZcOucR3lskAQJyfIEYQ06sF /ddsLPThlHL14CeALHqZujN2XD5jNXEvNzIjpcZ9UkBD/JCZYTl5haae2jxUexHWF/4i bFHw== X-Gm-Message-State: AOJu0Yz3zVC0Uksof2AruuGJtKqRMa95x7+akh4mT4lbm6Y1ZAnphf/q MduYka+1gbzSbCKSmb7H2Ojt9y82aLACQ/eWMVhIFm7DlBz2wQomKomJl/2mOhrlx/udbcZKo+a Tp4Lf X-Gm-Gg: AR+sD13d/mfEJ8MXejlz9T8i4T499CV3GQKQgbTnC48q3iR+6mp5oz0zBqoL1Es5/qj o2A2aLQ8cL6FtNKG5BKSuVfTJS3n7w+gYFlVrwb75L0OQ2YmpLkuTA/U6wDaq6AvVpNWUQqZ1RQ 17lCsi8UFjBlIyJN5Rra/l5yhCNatcTs7JZkTBVoXZI9SRMVqAQnP/jbuBQgv9let3yD6fxXiew ZlJNnInJEX0/cnY+LzSh8sDCbsqgdOYejixByGudgFF/KnMoTQ4LWylbm1VwdbeuxH1Grq8dNgB YyhWqAENJXvDezJ2kxQa8x65PDJ9UHCzEWN97hKr5A64JDKbBnkTxtUWWrrKt2vnVef3AMglkhW 7NItR3kTng182x8FLOY5ZUrPKca1Zyr7O/eRja1sHr3KsD+8Z1IYemVeeVZQ7EH+6iCmVoBrQUz DcCd8PZ7+FZnHnt5QuaKGkCw82QYc266JRurz7Ixq+90TqbF+W8gqMFBtZnl3Fctzs X-Received: by 2002:a05:6000:2586:b0:477:3986:9b54 with SMTP id ffacd0b85a97d-47fec62e2b4mr13722134f8f.20.1785961512724; Wed, 05 Aug 2026 13:25:12 -0700 (PDT) Received: from raptor ([62.108.198.223]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-47ff79b431dsm67029f8f.13.2026.08.05.13.25.11 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 05 Aug 2026 13:25:12 -0700 (PDT) From: Paul-Antoine Arras <parras@baylibre.com> To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, Paul-Antoine Arras <parras@baylibre.com> Subject: [PATCH 0/5] openmp: Support array-of-pointers (multi-segment) noncontiguous array sections and array-shaping casts Date: Wed, 5 Aug 2026 22:23:38 +0200 Message-ID: <20260805202343.2868178-1-parras@baylibre.com> X-Mailer: git-send-email 2.53.0 MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Spam-Status: No, score=-5.8 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP shortcircuit=no autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list <gcc-patches.gcc.gnu.org> List-Unsubscribe: <https://gcc.gnu.org/mailman/options/gcc-patches>, <mailto:gcc-patches-request@gcc.gnu.org?subject=unsubscribe> List-Archive: <https://gcc.gnu.org/pipermail/gcc-patches/> List-Post: <mailto:gcc-patches@gcc.gnu.org> List-Help: <mailto:gcc-patches-request@gcc.gnu.org?subject=help> List-Subscribe: <https://gcc.gnu.org/mailman/listinfo/gcc-patches>, <mailto:gcc-patches-request@gcc.gnu.org?subject=subscribe> Errors-To: gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org |
| Series |
openmp: Support array-of-pointers (multi-segment) noncontiguous array sections and array-shaping casts
|
|
Message
Paul-Antoine Arras
Aug. 5, 2026, 8:23 p.m. UTC
GCC already lets a "target"/"target update" map clause describe a
noncontiguous array section or an array-shaping cast over a single
array or a single pointer indirection, e.g.:
int a[N][M];
#pragma omp target update to(a[2:3:4][0:4:2])
int *p;
#pragma omp target update to((([N]) p)[2:3:4])
What wasn't supported is a section or shape-cast that crosses more
than one pointer indirection, e.g. an array-of-pointers:
int *a[N];
#pragma omp target update to(a[2:3:4][0:4:2])
Each further indirection needs its own, independently-strided
"segment" of dimensions in the noncontiguous-array descriptor that
gets built for the directive and walked on the target side, since the
memory reachable through each pointer in the outer dimension(s) is a
separate, unrelated allocation. This series adds that support, for C
and C++, on top of the existing single-segment (plain array section /
single pointer indirection) machinery [1].
Patch 1 adds the new GOMP_MAP_SHAPE_DIM map kind and the
OMP_CLAUSE_MAP_GRID_DIM_POINTER clause flag used by the rest of the
series to mark a dimension as selecting through a pointer.
Patch 2 teaches the C and C++ front ends to build the right clause
chain for these cases (array-of-pointers sections, array-shaping
casts crossing a pointer, decayed array parameters), with tests
scanning the "original" dump to check the clauses built for each case.
Patches 3 and 4 do the corresponding lowering work: patch 3 is a
pure, non-functional outlining of the (pre-existing, single-segment)
descriptor-construction code in lower_omp_target into its own
function, lower_omp_target_grid_desc; patch 4 generalizes that
function to multiple segments and the new clause kind, and adds the
"lower" dump scans (checking the generated descriptor) to the tests
added in patch 2.
Patch 5 extends the libgomp runtime descriptor and its gather/scatter
loops to walk multiple segments, with C and C++ runtime tests for one,
two, and three segments.
This series targets devel/omp/gcc-16.
[1] The strided patches have not landed in mainline yet. They are on
the devel/omp/gcc-16 branch in the commits:
2a5fe5c4b3f OpenMP: Support strided and shaped-array updates for C++
40617ea1249 OpenMP: Array shaping operator and strided "target update" for C
1b6460d9eb6 OpenMP: Noncontiguous "target update" for Fortran
Paul-Antoine Arras (5):
openmp: Add GOMP_MAP_SHAPE_DIM and grid-dim pointer marker
prerequisites
openmp: Support pointer indirection in array-shaping casts and array
sections
openmp: Outline lower_omp_target_grid_desc from lower_omp_target
openmp: Generalize lower_omp_target_grid_desc to segments and pointer
sections
openmp: Support multi-segment noncontiguous descriptors in libgomp
gcc/c-family/c-omp.cc | 151 +++-
gcc/c/c-parser.cc | 59 +-
gcc/c/c-typeck.cc | 144 +++-
gcc/cp/decl.cc | 50 +-
gcc/cp/parser.cc | 69 +-
gcc/cp/semantics.cc | 114 ++-
gcc/gimplify.cc | 18 +-
gcc/omp-general.cc | 5 +-
gcc/omp-low.cc | 798 +++++++++++-------
.../c-c++-common/gomp/array-section-1.c | 40 +
.../c-c++-common/gomp/array-section-2.c | 22 +
.../c-c++-common/gomp/array-section-3.c | 45 +
.../c-c++-common/gomp/array-section-4.c | 28 +
.../c-c++-common/gomp/array-section-5.c | 32 +
.../c-c++-common/gomp/array-section-6.c | 19 +
.../c-c++-common/gomp/array-section-8.c | 34 +
.../c-c++-common/gomp/array-section-9.c | 28 +
.../gomp/target-update-iterators-4.c | 19 +
gcc/testsuite/g++.dg/gomp/array-shaping-3.C | 71 ++
gcc/testsuite/g++.dg/gomp/array-shaping-4.C | 37 +
.../g++.dg/gomp/bad-array-shaping-9.C | 51 ++
gcc/testsuite/gcc.dg/gomp/array-shaping-9.c | 25 +
.../gcc.dg/gomp/bad-array-section-c-9.c | 60 ++
.../gcc.dg/gomp/bad-array-shaping-c-8.c | 45 +
gcc/tree-pretty-print.cc | 9 +
gcc/tree.h | 7 +
include/gomp-constants.h | 3 +-
libgomp/libgomp.h | 24 +-
libgomp/target.c | 176 +++-
.../testsuite/libgomp.c++/array-section-1.C | 204 +++++
.../testsuite/libgomp.c++/array-section-2.C | 130 +++
.../libgomp.c-c++-common/array-section-1.c | 181 ++++
.../libgomp.c-c++-common/array-section-10.c | 69 ++
.../libgomp.c-c++-common/array-section-11.c | 73 ++
.../libgomp.c-c++-common/array-section-2.c | 75 ++
.../libgomp.c-c++-common/array-section-3.c | 135 +++
.../libgomp.c-c++-common/array-section-4.c | 134 +++
.../libgomp.c-c++-common/array-section-5.c | 136 +++
.../libgomp.c-c++-common/array-section-6.c | 75 ++
.../libgomp.c-c++-common/array-section-7.c | 104 +++
.../libgomp.c-c++-common/array-section-8.c | 87 ++
.../libgomp.c-c++-common/array-section-9.c | 120 +++
42 files changed, 3213 insertions(+), 493 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-4.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-5.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-6.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-8.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-9.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
create mode 100644 gcc/testsuite/g++.dg/gomp/array-shaping-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/array-shaping-4.C
create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C
create mode 100644 gcc/testsuite/gcc.dg/gomp/array-shaping-9.c
create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c
create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c
create mode 100644 libgomp/testsuite/libgomp.c++/array-section-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/array-section-2.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-10.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-11.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-3.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-4.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-5.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-6.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-7.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-8.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-section-9.c