From patchwork Wed Aug 5 20:23:39 2026 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 140681 Return-Path: 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 9AD4A4BAE7E3 for ; Wed, 5 Aug 2026 20:26:58 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 9AD4A4BAE7E3 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=m7Ke28GU X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x32b.google.com (mail-wm1-x32b.google.com [IPv6:2a00:1450:4864:20::32b]) by sourceware.org (Postfix) with ESMTPS id 907284BAE7E8 for ; Wed, 5 Aug 2026 20:25:39 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 907284BAE7E8 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 907284BAE7E8 Authentication-Results: sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32b ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961539; cv=none; b=tI2gl1AAnRiNzd4C540QcAE6tStGvs+/7lKo+C509tKlzpphQy1+mB/U0Mu+2xLMy25rF8OuKI5gCydvdCx3OtnetqNkqdrdKjCO/OpAsPLic+EyBxhq7MWqYNSob7/IaAYBjJ7K0X0xTTfVITCQNXlLCnXkMMWrb63UxrMCNpc= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961539; c=relaxed/simple; bh=EwtZs7JB7Adpzcml7fLCf8PlAqtxS09b0xlW/+TszKo=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=ERiISW52AfHwK4Vfyio4xkbkZkPwoFcsp/9NFu1odKJtDKVYghurSdVbaqMzVcNoavvwIq1A+OlKUjMz6AQJVnrKexBih+5KqqME/HOZ+A6ccgsn09Fl0uFKPbCy8SaL9q76gux6h3Fb54ag3w089M7SC3KTAAxnkTpkJOlUo0A= 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=m7Ke28GU DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 907284BAE7E8 Received: by mail-wm1-x32b.google.com with SMTP id 5b1f17b1804b1-4954aff6088so15020775e9.3 for ; Wed, 05 Aug 2026 13:25:39 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1785961538; x=1786566338; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to:content-type; bh=8MWDrBEWI4UbQKTjvE/hTMrwbI/d/ICITaOqfzLL+Sw=; b=m7Ke28GUdMr9EfJiI/+yzpS0HK40h3iGxckajAWmrQFu69zez1mUv2scvGC6tXZIJ4 kbXb+sTUqML2ZQUkK9/8jbjJLX1K92jE4jtJbz9TxOUqrZ1DWlLO61jUvh61TBpMsiKt meiDGlZ69fRE+NWCtJEzOMTxNu4FqT03vVckako0XSWhsCBpg7euqFJ09MYgjKqSxiS6 eT+QBQlSH6BsokLx2o1HiTCRqPRKFigj/CHYpHYIHooOI+rsPuljDJkKOfpXQYh5mK9v sd6ICrLGHgYiOWhXXsTschStMHHVzkn4DxZlh4nxyhKR3YErxZCrbThG1sCgwpGxNm9I 3sPA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1785961538; x=1786566338; h=content-transfer-encoding:mime-version:references:in-reply-to :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=8MWDrBEWI4UbQKTjvE/hTMrwbI/d/ICITaOqfzLL+Sw=; b=msdsbWq0vXYL8MbcLJ1AYQ9e3eawLUKHvV8Prw3YCw886sKop9mWjYOBLGNkyRcThi 8VRHYUz8RQlwcvoFdf8pfbxTU1WgP47N3B127olrtuBpwxCfv6hQC5EME5OQQf3O/Acn MDB24oR0YUPANkMrngAyPkAZJikr6ZDp2VhFU6oDp31gWpKjIzSzhBFiu/d7MT4xj3jG Pq91oU8iYM3wkCqXPB6tLsh5p1TlSCHDVDu6B1/hbrAeVHqYy4rOYUYSYGrxGvEcfeUI lz9DSJdoSUXyRjV8MUAz3H864G5z4+SaYZVYBTgO72Zzpx9nNZ+s9hS6mC9RYhtFv0v2 5xyA== X-Gm-Message-State: AOJu0YxhZsCTXY6YF2dbbgFeYGqMvF2MaCYJzyD1rw6r1DiXic1a9BU8 kRPfKV54pyH2K9tATEDP1owQ77jLNFuQ8N5uFbIXvQyDBOWC0FLa9dn4SZIxNQgz/Mrwd+AH/za rtB9q X-Gm-Gg: AR+sD10Z3Rlftn9p640qwIRXg0jzldl4vFm2EYa7cpe/02Wh+8tyz0FWnXtJqLIjKvO xX2MgfB/ZaAAybwMka3MvoHYSt0vhYP5X/Nmha65tekTHeOek0z5DiqAOmB6BHnCg+Wg1R/lr8d e1qRfmuOEY6AVcmUi/IW8ed6pxQ04RNVsYZPn9pnEBRe5cdZEcos2ABlxC8h5OGy9tvM/X5ldIl P6fijmeRKf8+aMMfz/4q7S8DH074L1jjjlSpyKqJ3ZEJ8qKbwAU7hlNB3dvKxLvyG04R5Bykcau UwG7gF2AxXJsp8G7lWDAJHEs8NKv+iIEekUYumQAFecC5MxEE1Yf7TDIxQKdqUwFAOMiCZQD940 8eEnkNZVGDXaz7uIeqIR58KxtRy61rRkKkGY9rT3t9DDX624gVckfnrQaYJBNwXN3oyntPzhPaI 4umqnTxDqyoHETNyqipr8Dc5q0TPJpmnk4qQITdabHpYLoOsw7tuEbyQ== X-Received: by 2002:a05:600c:4688:b0:495:6bc9:62b0 with SMTP id 5b1f17b1804b1-4994e7d1487mr122013555e9.17.1785961537940; Wed, 05 Aug 2026 13:25:37 -0700 (PDT) Received: from raptor ([62.108.198.223]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-47ff79b431dsm67029f8f.13.2026.08.05.13.25.37 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 05 Aug 2026 13:25:37 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, Paul-Antoine Arras Subject: [PATCH 1/5] openmp: Add GOMP_MAP_SHAPE_DIM and grid-dim pointer marker prerequisites Date: Wed, 5 Aug 2026 22:23:39 +0200 Message-ID: <20260805202343.2868178-2-parras@baylibre.com> X-Mailer: git-send-email 2.53.0 In-Reply-To: <20260805202343.2868178-1-parras@baylibre.com> References: <20260805202343.2868178-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, 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 List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org Introduce the new GOMP_MAP_SHAPE_DIM map kind, the OMP_CLAUSE_MAP_GRID_DIM_POINTER accessor, and their tree-pretty-print support ahead of the front-end and lowering changes that use them. gcc/ChangeLog: * tree-pretty-print.cc (dump_omp_clause): Handle GOMP_MAP_SHAPE_DIM and print [pointer]/[pointer set] annotations. * tree.h (OMP_CLAUSE_MAP_GRID_DIM_POINTER): New. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_SHAPE_DIM. --- gcc/tree-pretty-print.cc | 9 +++++++++ gcc/tree.h | 7 +++++++ include/gomp-constants.h | 3 ++- 3 files changed, 18 insertions(+), 1 deletion(-) diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 0959397a149..71d430d7330 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1167,6 +1167,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_GRID_STRIDE: pp_string (pp, "grid_stride"); break; + case GOMP_MAP_SHAPE_DIM: + pp_string (pp, "shape_dim"); + break; case GOMP_MAP_UNSET: pp_string (pp, "unset"); break; @@ -1273,6 +1276,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, " [runtime_implicit]"); if (OMP_CLAUSE_MAP_GIMPLE_ONLY (clause)) pp_string (pp, " [gimple only]"); + if (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET + && !OMP_CLAUSE_SIZE (clause)) + pp_string (pp, " [pointer set]"); + if (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_GRID_DIM + && OMP_CLAUSE_MAP_GRID_DIM_POINTER (clause)) + pp_string (pp, " [pointer]"); } pp_right_paren (pp); break; diff --git a/gcc/tree.h b/gcc/tree.h index 6aa263aaada..72306428920 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1971,6 +1971,13 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_POINTS_TO_READONLY(NODE) \ TREE_CONSTANT (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) +/* Nonzero on a GOMP_MAP_GRID_DIM clause if the dimension it describes + selects into a pointer (requiring a dereference to reach the next + dimension) rather than a regular array dimension. */ +#define OMP_CLAUSE_MAP_GRID_DIM_POINTER(NODE) \ + TREE_THIS_VOLATILE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + + /* Same as above, for use in OpenACC cache directives. */ #define OMP_CLAUSE__CACHE__READONLY(NODE) \ TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 96ff5ee4b44..7bf71f9d55d 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -252,7 +252,8 @@ enum gomp_map_kind definition (only for Fortran at present). */ GOMP_MAP_MAPPING_GROUP = (GOMP_MAP_LAST | 11), GOMP_MAP_GRID_DIM = (GOMP_MAP_LAST | 12), - GOMP_MAP_GRID_STRIDE = (GOMP_MAP_LAST | 13) + GOMP_MAP_GRID_STRIDE = (GOMP_MAP_LAST | 13), + GOMP_MAP_SHAPE_DIM = (GOMP_MAP_LAST | 14) }; #define GOMP_MAP_COPY_TO_P(X) \ From patchwork Wed Aug 5 20:23:40 2026 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 140680 Return-Path: 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 A7A344BB1C3F for ; Wed, 5 Aug 2026 20:26:55 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A7A344BB1C3F 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=hxGSdk7e 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 DB1C24BAE7E6 for ; Wed, 5 Aug 2026 20:25:40 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org DB1C24BAE7E6 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 DB1C24BAE7E6 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=1785961541; cv=none; b=D+lEsQbRaKS7KBzAObhr9a/I2lrYquxDJKVWGzVtXsiJJd8Dd1YA4QmTjCWt+42WtYeWGtTSi9K9yzn/icsq+ZTub7UEi/yws6Kb5nP/MFBxTzNxCbmxMe1IUT2ZbzWhkqHD9ruAmQwTGJOOYhVzuXHcI9zZw4nyQt1qQ7wkF/8= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961541; c=relaxed/simple; bh=kLkPszk9z/scZ7Kddd6CFSrq3dtuq30cDygj3tmVyog=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=g9e3iSh7RRGROCSxC4w7jHEJuwpzxpAIdD1JaqyIHMgIAiN9Tq6mzlMCissVppkJkwBcIirwAO9DmWGK3/FqvUeSOPZerPNZmmZc7YXUBtbbw45jVGiyoyAi5WT2Rpsw8Ovbv280BPOVSseLH/WtAUR7Jry7qzXDwNHsQFBvVYA= 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=hxGSdk7e DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org DB1C24BAE7E6 Received: by mail-wr1-x42e.google.com with SMTP id ffacd0b85a97d-472326ca506so1030322f8f.2 for ; Wed, 05 Aug 2026 13:25:40 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1785961540; x=1786566340; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to:content-type; bh=BDJWSPXV9Sv0aYibjKsPpaeoAitjnxYYjfVj4a6uDzQ=; b=hxGSdk7emBJ3aNlqQdrmptea9OBmKN0xD6FkmflSZdEUb41srf4Sd/e6qKW1HJVZGa Mo/iuEehxRHFevqAPAeI7CpTXtVUysAvYyaP3s15fuEvurtXYQ+4iD7snUjaBJUhSG+v XR7+fqCQRYJiyx+tiKKAPI8IdSAlN+Sdo+GyFcNqsHhcf78/L43g4rv1QcNGl11f44gd gDoD4cN1XVTKkRVjINPbGFnJGO5D/3s4/gzqKZpfHZpICoEszByYRU56WkFYAX8GIGFs TeCDxh5YPevYXGrCXUTezgnYBy1bo1Ck5YUJNQSmacZsGBmu/U1dzA5k4H63Fhwmmz7k WzIw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1785961540; x=1786566340; h=content-transfer-encoding:mime-version:references:in-reply-to :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=BDJWSPXV9Sv0aYibjKsPpaeoAitjnxYYjfVj4a6uDzQ=; b=SUPnVl/4IwesVlEN3sKYN5CrNch9bQzWE5j/dd5t33lp+RQSw4fXUDC/AjZN9sEP01 LISjS9iGcn95STJhw+lwdI9eu4yrErfZVdpoKftE18yhsZBODsXCJQC+B9zYUWQ0zM3C FTYv+E6ZbMb+137xlo/reUSpaHP9jovAo8OWrDQemJffENjcHT92CstxzC51aBRGiabl 58ZJahYjsiW3fprYkpa2BeW5fuc0AiM8LJGDp9ePikJNWq+g0jxUuZbyTJF7mDYDYSZQ kPc2Q+aNKf1KqV+gWhwj3QuYo8Gg79a4n0Vj0i3ebyrTrOLsiDFzjH3UsuQQGjborU+N klPw== X-Gm-Message-State: AOJu0YwKe6/hcmxVMirFmWmUtLwZROv7TThAs3q9h818s6a/FohEj65n lVR0xXLqSQnOXhIB1Bz6XDXBx7ysUKm5hbDkW5dWy+OncOoQJt0OSa3lR1NHJY5qsr6Tl2zXMLj siYsA X-Gm-Gg: AR+sD130+V+pCOaiKlyfRYD2LcAsaw5hn6plXVHxq6Xr6lIce2Lj/+ve5xMCvu+Veiz sRsG3xWrhVI0pXl14cFG1JAfEWO6ZE0PLDA2iwfyZ3vaMpXRvSBeb4okhxRAKgBCUL9U2ItQ1AX 7yR7BpNN8WWfZOl+Se6yuPbccY0cFdvtj3KfWKwzLVEhmIYBZLo2yKknVPWPytu7A3YGWmJbxje +BK2/XyNAuKqoer6yye8hJPKXrooqmS7JVIXe9V7FmqJO9RWHad0qhOFxZRReZ/Sd88Vq/3CxSU weBD72OgpOItZo12goTjPeow4tuAEW0AVTTLOWkzlFb4xMVdhmoOxqxJTY/5uXauWm60wrgOnkj P56o4OBO6xHMH0pw5ZcBZJ6Sd1SlD4SVNv2vzihrCnFnO+6MbLbWbF7v2L57WLn22epDVUWb/Q+ J/CaDpT34xIeU/+6aorjF6wUeOygIkAcNfWOVsisHYcGv03hiQme5ocw== X-Received: by 2002:a05:6000:2c0d:b0:47f:6b9a:9d54 with SMTP id ffacd0b85a97d-47fec4ebd82mr16348706f8f.7.1785961539210; Wed, 05 Aug 2026 13:25:39 -0700 (PDT) Received: from raptor ([62.108.198.223]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-47ff79b431dsm67029f8f.13.2026.08.05.13.25.38 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 05 Aug 2026 13:25:38 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, Paul-Antoine Arras Subject: [PATCH 2/5] openmp: Support pointer indirection in array-shaping casts and array sections Date: Wed, 5 Aug 2026 22:23:40 +0200 Message-ID: <20260805202343.2868178-3-parras@baylibre.com> X-Mailer: git-send-email 2.53.0 In-Reply-To: <20260805202343.2868178-1-parras@baylibre.com> References: <20260805202343.2868178-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.9 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, PROLO_LEO1, 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 List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org Array-shaping casts and noncontiguous array sections were already partially supported; what is new here is handling further pointer indirection within them, i.e. array-of-pointers sections and array-shaping casts that cross more than one pointer. Each further indirection starts a new "segment" of dimensions, recorded by marking the GOMP_MAP_GRID_DIM clause for that dimension with OMP_CLAUSE_MAP_GRID_DIM_POINTER; a GOMP_MAP_SHAPE_DIM clause is emitted alongside dimensions introduced by an array-shaping cast so the lowering pass can size them. This also fixes crashes/leaks uncovered along the way (member access via a dependent object, 2+-dimensional array-of-pointers in templates), rejects array-shaping casts over a range of pointers per the OpenMP spec, accepts decayed array parameters, and teaches omp_parse_noncontiguous_array to look through the VIEW_CONVERT_EXPR an array-shaping cast introduces over an OMP_ARRAY_SECTION. Add tests scanning the "original" dump to check the clause chains built for each case; the corresponding "lower" dump checks for the generated descriptors are added once the omp-low.cc lowering support lands. gcc/c-family/ChangeLog: * c-omp.cc (omp_expand_grid_dim): Support array-of-pointers sections spanning multiple segments and dimensions introduced by an array-shaping cast, marking dimensions that select through a pointer and emitting GOMP_MAP_SHAPE_DIM clauses for shape-cast dimensions. (omp_handle_noncontig_array): Adjust to omp_expand_grid_dim's new signature and synthesize whole-span GOMP_MAP_GRID_DIM/ GOMP_MAP_SHAPE_DIM clauses for any shape-cast dimensions left over. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_variable_list): Accept array-shaping casts over decayed array parameters. * c-typeck.cc (create_omp_arrayshape_type): Reject array-shaping casts over a range of pointers. (handle_omp_array_sections_1): Support array-of-pointers sections. (handle_omp_array_sections): Likewise. (c_finish_omp_clauses): Likewise. gcc/cp/ChangeLog: * decl.cc (cp_omp_create_arrayshape_type): Reject array-shaping casts over a range of pointers. * parser.cc (cp_parser_omp_var_list_no_open): Accept array-shaping casts over decayed array parameters. * semantics.cc (handle_omp_array_sections_1): Support array-of-pointers sections; fix crash on array-shaping cast of a member via a dependent object and leaks for 2+-dimensional array-of-pointers in templates. (finish_omp_clauses): Support array-of-pointers sections. gcc/ChangeLog: * omp-general.cc (omp_parse_noncontiguous_array): Look through the VIEW_CONVERT_EXPR an array-shaping cast introduces over an OMP_ARRAY_SECTION. gcc/testsuite/ChangeLog: * c-c++-common/gomp/array-section-1.c: New test. * c-c++-common/gomp/array-section-2.c: New test. * c-c++-common/gomp/array-section-3.c: New test. * c-c++-common/gomp/array-section-4.c: New test. * c-c++-common/gomp/array-section-5.c: New test. * c-c++-common/gomp/array-section-6.c: New test. * c-c++-common/gomp/array-section-8.c: New test. * c-c++-common/gomp/array-section-9.c: New test. * c-c++-common/gomp/target-update-iterators-4.c: New test. * g++.dg/gomp/array-shaping-3.C: New test. * g++.dg/gomp/array-shaping-4.C: New test. * g++.dg/gomp/bad-array-shaping-9.C: New test. * gcc.dg/gomp/array-shaping-9.c: New test. * gcc.dg/gomp/bad-array-section-c-9.c: New test. * gcc.dg/gomp/bad-array-shaping-c-8.c: New test. --- 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/omp-general.cc | 5 +- .../c-c++-common/gomp/array-section-1.c | 33 ++++ .../c-c++-common/gomp/array-section-2.c | 15 ++ .../c-c++-common/gomp/array-section-3.c | 38 +++++ .../c-c++-common/gomp/array-section-4.c | 21 +++ .../c-c++-common/gomp/array-section-5.c | 25 +++ .../c-c++-common/gomp/array-section-6.c | 18 +++ .../c-c++-common/gomp/array-section-8.c | 27 ++++ .../c-c++-common/gomp/array-section-9.c | 21 +++ .../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 ++++++ 22 files changed, 973 insertions(+), 125 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 diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 01721cafad0..c4941166206 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3718,60 +3718,110 @@ omp_expand_access_chain (tree *pc, tree expr, return pc; } -static tree * -omp_expand_grid_dim (location_t loc, tree *pc, tree decl) -{ - if (TREE_CODE (decl) == OMP_ARRAY_SECTION) - pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0)); - else - return pc; +/* Append a GOMP_MAP_GRID_DIM (and possibly GOMP_MAP_GRID_STRIDE and + GOMP_MAP_SHAPE_DIM) clause for each dimension of the OMP_ARRAY_SECTION chain + in DECL onto the clause chain at *PC, base to outer. PC points to the + insertion point in the clause chain. DECL is the OMP_ARRAY_SECTION being + expanded, or the base declaration once the chain is exhausted. *TYPE is set + to the base declaration's type, then peeled through recursion so a dimension + is marked OMP_CLAUSE_MAP_GRID_DIM_POINTER when it is a pointer rather than an + array (output only). *FIRST is true only for the first dimension, so that it + is never marked as pointer even when the base type has decayed (output only). + *SHAPE_TYPE is set to the array-shaping type for the GOMP_MAP_SHAPE_DIM + clause (output only). Returns the updated chain insertion point. */ - tree c = *pc; +static tree * +omp_expand_grid_dim (location_t loc, tree *pc, tree decl, tree *type, + bool *first, tree *shape_type) +{ + if (TREE_CODE (decl) == OMP_ARRAY_SECTION || TREE_CODE (decl) == ARRAY_REF) + { + tree op = TREE_OPERAND (decl, 0); + if (TREE_CODE (op) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_OPERAND (op, 0)) == OMP_ARRAY_SECTION) + { + /* Array section within an array-shaping cast. */ + pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (op, 0), type, + first, shape_type); + *shape_type = TREE_TYPE (op); + } + else + pc = omp_expand_grid_dim (loc, pc, op, type, first, shape_type); + } + else + { + *type = TREE_TYPE (decl); + *first = true; + *shape_type = NULL_TREE; + return pc; + } + + tree nc = OMP_CLAUSE_CHAIN (*pc); tree low_bound = TREE_OPERAND (decl, 1); tree length = TREE_OPERAND (decl, 2); tree stride = TREE_OPERAND (decl, 3); - tree cd = build_omp_clause (loc, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (cd, GOMP_MAP_GRID_DIM); - OMP_CLAUSE_DECL (cd) = unshare_expr (low_bound); - OMP_CLAUSE_SIZE (cd) = unshare_expr (length); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_GRID_DIM); + OMP_CLAUSE_DECL (c) = unshare_expr (low_bound); + OMP_CLAUSE_SIZE (c) = length ? unshare_expr (length) : size_one_node; + if (!*first && TREE_CODE (*type) == POINTER_TYPE) + OMP_CLAUSE_MAP_GRID_DIM_POINTER (c) = 1; + OMP_CLAUSE_CHAIN (*pc) = c; + pc = &OMP_CLAUSE_CHAIN (*pc); if (stride && !integer_onep (stride)) { - tree cs = build_omp_clause (loc, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (cs, GOMP_MAP_GRID_STRIDE); - OMP_CLAUSE_DECL (cs) = unshare_expr (stride); - - OMP_CLAUSE_CHAIN (cs) = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (cd) = cs; - OMP_CLAUSE_CHAIN (c) = cd; - pc = &OMP_CLAUSE_CHAIN (cd); + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_GRID_STRIDE); + OMP_CLAUSE_DECL (c) = unshare_expr (stride); + OMP_CLAUSE_CHAIN (*pc) = c; + pc = &OMP_CLAUSE_CHAIN (*pc); } - else + if (*shape_type) { - OMP_CLAUSE_CHAIN (cd) = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (c) = cd; - pc = &OMP_CLAUSE_CHAIN (c); - } + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_SHAPE_DIM); + tree dtype = TYPE_DOMAIN (*shape_type); + tree minval = TYPE_MIN_VALUE (dtype); + tree maxval = TYPE_MAX_VALUE (dtype); + minval = fold_convert (sizetype, minval); + maxval = fold_convert (sizetype, maxval); + tree dim = size_binop (MINUS_EXPR, maxval, minval); + dim = size_binop (PLUS_EXPR, dim, size_one_node); + OMP_CLAUSE_DECL (c) = dim; + OMP_CLAUSE_CHAIN (*pc) = c; + pc = &OMP_CLAUSE_CHAIN (*pc); + tree elt = TREE_TYPE (*shape_type); + *shape_type = TREE_CODE (elt) == ARRAY_TYPE ? elt : NULL_TREE; + } + OMP_CLAUSE_CHAIN (c) = nc; + + *first = false; + *type = *shape_type ? *shape_type : TREE_TYPE (*type); return pc; } +/* Replace clause C, a TO/FROM clause on a non-contiguous array section, at + *PC with a GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause describing BASE + (the section's base declaration), followed by the GOMP_MAP_GRID_DIM/ + GOMP_MAP_GRID_STRIDE clauses for each dimension of C's OMP_ARRAY_SECTION + chain. Returns the updated chain insertion point. */ + tree * omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base) { - tree type; + tree eltype = TREE_TYPE (base); - if (POINTER_TYPE_P (TREE_TYPE (base))) - type = TREE_TYPE (TREE_TYPE (base)); - else - type = strip_array_types (TREE_TYPE (base)); + while (TREE_CODE (eltype) == ARRAY_TYPE || POINTER_TYPE_P (eltype)) + eltype = TREE_TYPE (eltype); tree c_map = build_omp_clause (loc, OMP_CLAUSE_MAP); OMP_CLAUSE_DECL (c_map) = unshare_expr (base); /* Use the element size (or pointed-to type size) here. */ - OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (type); + OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (eltype); switch (OMP_CLAUSE_CODE (c)) { @@ -3789,7 +3839,44 @@ omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base) *pc = c_map; - return omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c)); + tree dtype; + bool first; + tree shape_type; + pc = omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c), &dtype, &first, + &shape_type); + + /* Build whole-span dimensions for remaining shape_type layers. */ + tree nc = OMP_CLAUSE_CHAIN (*pc); + tree dc = NULL_TREE; + while (shape_type != NULL_TREE && TREE_CODE (shape_type) == ARRAY_TYPE) + { + /* Compute whole-span length. */ + tree dtype = TYPE_DOMAIN (shape_type); + gcc_assert (integer_zerop (TYPE_MIN_VALUE (dtype))); + tree length = TYPE_MAX_VALUE (dtype); + length = size_binop (PLUS_EXPR, length, size_one_node); + + /* Build grid_dim node. */ + dc = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (dc, GOMP_MAP_GRID_DIM); + OMP_CLAUSE_DECL (dc) = size_zero_node; + OMP_CLAUSE_SIZE (dc) = length; + OMP_CLAUSE_CHAIN (*pc) = dc; + pc = &OMP_CLAUSE_CHAIN (*pc); + + /* Build shape_dim node. */ + dc = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (dc, GOMP_MAP_SHAPE_DIM); + OMP_CLAUSE_DECL (dc) = length; + OMP_CLAUSE_CHAIN (*pc) = dc; + pc = &OMP_CLAUSE_CHAIN (*pc); + + shape_type = TREE_TYPE (shape_type); + } + if (dc != NULL_TREE) + OMP_CLAUSE_CHAIN (dc) = nc; + + return pc; } /* Translate "array_base_decl access_method" to OMP mapping clauses. */ diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 1c54af9910e..3d5a4825a10 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -17017,6 +17017,18 @@ c_parser_omp_variable_list (c_parser *parser, sections--; } + /* The previous loop may have uncovered a shaping cast. DECL_P + excludes such a cast over a further array section (reshaping + e.g. "x[1][1:2]"): that case is left for omp_expand_grid_dim's + SHAPE_TYPE handling in c-omp.cc. */ + if (!reshaped_to && TREE_CODE (decl) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE + && DECL_P (TREE_OPERAND (decl, 0))) + { + reshaped_to = TREE_TYPE (decl); + decl = TREE_OPERAND (decl, 0); + } + /* The handling of INDIRECT_REF here in the presence of array-shaping operations is a little tricky. We need to avoid treating a pointer dereference as a unit-sized array @@ -17087,7 +17099,52 @@ c_parser_omp_variable_list (c_parser *parser, } } - if (reshaped_to) + if (reshaped_to && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + { + unsigned array_layers = 0; + tree basetype = TREE_TYPE (decl); + while (TREE_CODE (basetype) == ARRAY_TYPE) + { + basetype = TREE_TYPE (basetype); + array_layers++; + } + + /* A further level of indirection here (e.g. "int **x[N]") + is not supported (for now); treating the second pointer's + bytes as the reshaped array's elements would silently read + the wrong data. (See also + gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c). */ + if (TREE_CODE (TREE_TYPE (basetype)) == POINTER_TYPE) + { + sorry_at (loc, "more than one level of pointer " + "indirection in a noncontiguous target " + "update"); + decl = error_mark_node; + } + else if (dims.length () != array_layers) + { + error_at (loc, "too many array section specifiers " + "for %qT", reshaped_to); + decl = error_mark_node; + } + else + { + for (unsigned i = 0; i < array_layers; i++) + { + omp_dim &d = dims[dims.length () - 1 - i]; + decl = build_omp_array_section (loc, decl, + d.low_bound, d.length, + d.stride); + } + dims.truncate (0); + decl = build1_loc (loc, VIEW_CONVERT_EXPR, reshaped_to, + decl); + tree elems = c_array_type_nelts_total (reshaped_to); + decl = build_omp_array_section (loc, decl, size_zero_node, + elems, NULL_TREE); + } + } + else if (reshaped_to) { unsigned reshaped_dims = 0; diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 0c141520501..ff061044ee4 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -3877,13 +3877,35 @@ create_omp_arrayshape_type (tree expr, vec *omp_shape_dims) while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION || TREE_CODE (strip_sections) == ARRAY_REF) - strip_sections = TREE_OPERAND (strip_sections, 0); + { + /* The array-shaping operator only applies to a single pointer, not an + array section. */ + if (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION) + { + tree array = TREE_OPERAND (strip_sections, 0); + tree length = TREE_OPERAND (strip_sections, 2); + tree atype = TREE_TYPE (array); + tree eltype = atype != NULL_TREE ? TREE_TYPE (atype) : NULL_TREE; + if (eltype != NULL_TREE + && POINTER_TYPE_P (eltype) + && (length == NULL_TREE || !integer_onep (length))) + { + error ("OpenMP array shaping operator with non-pointer " + "argument"); + return error_mark_node; + } + } + strip_sections = TREE_OPERAND (strip_sections, 0); + } tree type = TREE_TYPE (strip_sections); if (TREE_CODE (type) == REFERENCE_TYPE) type = TREE_TYPE (type); + while (TREE_CODE (type) == ARRAY_TYPE) + type = TREE_TYPE (type); + if (TREE_CODE (type) != POINTER_TYPE) { error ("OpenMP array shaping operator with non-pointer argument"); @@ -15864,14 +15886,16 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above case though, as some lengths could be zero. NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array - section. */ + section. TYPE carries the type of the inner section through recursion + (output only). */ static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, - bool &non_contiguous, enum c_omp_region_type ort, int *discontiguous) + bool &non_contiguous, enum c_omp_region_type ort, + int *discontiguous, tree *type) { - tree ret, low_bound, length, stride, type; + tree ret, low_bound, length, stride; bool openacc = (ort & C_ORT_ACC) != 0; if (TREE_CODE (t) != OMP_ARRAY_SECTION) { @@ -15941,20 +15965,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, ret = convert_lvalue_to_rvalue (OMP_CLAUSE_LOCATION (c), ret, false, false); } + *type = TREE_TYPE (ret); return ret; } ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types, maybe_zero_len, first_non_one, - non_contiguous, ort, - discontiguous); + non_contiguous, ort, discontiguous, type); if (ret == error_mark_node || ret == NULL_TREE) return ret; - if (TREE_CODE (ret) == OMP_ARRAY_SECTION) - type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0))); - else - type = TREE_TYPE (ret); low_bound = TREE_OPERAND (t, 1); length = TREE_OPERAND (t, 2); stride = TREE_OPERAND (t, 3); @@ -16041,11 +16061,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && (TREE_CODE (length) != INTEGER_CST || integer_onep (length))) first_non_one++; } - if (TREE_CODE (type) == ARRAY_TYPE) + if (TREE_CODE (*type) == ARRAY_TYPE) { if (length == NULL_TREE - && (TYPE_DOMAIN (type) == NULL_TREE - || TYPE_MAX_VALUE (TYPE_DOMAIN (type)) == NULL_TREE)) + && (TYPE_DOMAIN (*type) == NULL_TREE + || TYPE_MAX_VALUE (TYPE_DOMAIN (*type)) == NULL_TREE)) { error_at (OMP_CLAUSE_LOCATION (c), "for unknown bound array type length expression must " @@ -16069,13 +16089,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } - if (TYPE_DOMAIN (type) - && TYPE_MAX_VALUE (TYPE_DOMAIN (type)) - && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (type))) - == INTEGER_CST) + if (TYPE_DOMAIN (*type) && TYPE_MAX_VALUE (TYPE_DOMAIN (*type)) + && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (*type))) == INTEGER_CST) { tree size - = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (type))); + = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (*type))); size = size_binop (PLUS_EXPR, size, size_one_node); if (TREE_CODE (low_bound) == INTEGER_CST) { @@ -16102,11 +16120,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } maybe_zero_len = true; } - else if (length == NULL_TREE - && first_non_one == types.length () - && tree_int_cst_equal - (TYPE_MAX_VALUE (TYPE_DOMAIN (type)), - low_bound)) + else if (length == NULL_TREE && first_non_one == types.length () + && tree_int_cst_equal ( + TYPE_MAX_VALUE (TYPE_DOMAIN (*type)), low_bound)) first_non_one++; } else if (length == NULL_TREE) @@ -16188,7 +16204,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } } } - else if (TREE_CODE (type) == POINTER_TYPE) + else if (TREE_CODE (*type) == POINTER_TYPE) { if (length == NULL_TREE) { @@ -16267,7 +16283,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) - types.safe_push (type); + types.safe_push (*type); /* We will need to evaluate lb more than once. */ tree lb = save_expr (low_bound); if (lb != low_bound) @@ -16284,6 +16300,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, length, stride); else ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound); + + *type = TREE_TYPE (*type); return ret; } @@ -16322,9 +16340,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY) && OMP_ITERATOR_DECL_P (*tp)) tp = &TREE_VALUE (*tp); - tree first = handle_omp_array_sections_1 (c, *tp, types, - maybe_zero_len, first_non_one, - non_contiguous, ort, discontiguous); + tree type; + tree first + = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, + non_contiguous, ort, discontiguous, &type); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -16343,8 +16362,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, if (tem == NULL_TREE) tem = TREE_VALUE (t); else - tem = build2 (COMPOUND_EXPR, TREE_TYPE (tem), - TREE_VALUE (t), tem); + tem = build2 (COMPOUND_EXPR, TREE_TYPE (tem), TREE_VALUE (t), + tem); } t = TREE_CHAIN (t); } @@ -16364,6 +16383,11 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, maybe_zero_len = true; bool higher_discontiguous = false; + /* Whether the first dimension uses the extended array section syntax, + i.e. whether it specifies a length or a stride. */ + bool first_extended = false; + /* Whether any dimension container except the last is a pointer. */ + bool non_last_pointer = false; for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; t = TREE_OPERAND (t, 0)) @@ -16427,10 +16451,15 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, full_span = true; } + tree op = TREE_OPERAND (t, 0); + bool shaped + = TREE_CODE (op) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (op)) == ARRAY_TYPE + && TREE_CODE (TREE_OPERAND (op, 0)) == OMP_ARRAY_SECTION; if (!integer_onep (stride) || (higher_discontiguous - && (!integer_zerop (low_bound) - || !full_span))) + && (!integer_zerop (low_bound) || !full_span)) + || shaped) *discontiguous = 2; if (!integer_onep (stride) @@ -16440,20 +16469,29 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, } if (!maybe_zero_len && i > first_non_one) { + bool full_span; if (integer_nonzerop (low_bound)) goto is_noncontiguous; - if (length != NULL_TREE - && TREE_CODE (length) == INTEGER_CST - && TYPE_DOMAIN (types[i]) - && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])) - && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))) - == INTEGER_CST) + if (length != NULL_TREE && TREE_CODE (length) == INTEGER_CST) { - tree size; - size = size_binop (PLUS_EXPR, - TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])), - size_one_node); - if (!tree_int_cst_equal (length, size)) + /* A pointer-typed dimension has no static bound to compare + against, so it can never be proven to span the whole + dimension; treat it conservatively as not full span. */ + full_span = false; + if (TREE_CODE (types[i]) == ARRAY_TYPE + && TYPE_DOMAIN (types[i]) + && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])) + && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))) + == INTEGER_CST) + { + tree size; + size + = size_binop (PLUS_EXPR, + TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])), + size_one_node); + full_span = tree_int_cst_equal (length, size); + } + if (!full_span) { is_noncontiguous: if (discontiguous && *discontiguous) @@ -16538,7 +16576,24 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, else size = size_binop (MULT_EXPR, size, l); } + + if (i == 0 + && ((length != NULL_TREE && !integer_onep (length)) + || !integer_onep (stride))) + first_extended = true; + if (i < num - 1 && TREE_CODE (types[i]) == POINTER_TYPE + && /* accept decayed array */ !( + i == 0 && TREE_CODE (TREE_TYPE (types[i])) == ARRAY_TYPE)) + non_last_pointer = true; } + + if (first_extended && non_last_pointer) + error_at ( + OMP_CLAUSE_LOCATION (c), + "only the last dimension can refer to a pointer when the first " + "dimension uses the extended array section syntax in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + if (non_contiguous) { int kind = OMP_CLAUSE_MAP_KIND (c); @@ -17942,7 +17997,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_SHAPE_DIM) break; /* FALLTHRU */ case OMP_CLAUSE_TO: diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc index 1d67999cfef..cb590165091 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -13645,17 +13645,60 @@ cp_omp_create_arrayshape_type (location_t loc, tree expr, while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION || TREE_CODE (strip_sections) == ARRAY_REF) - strip_sections = TREE_OPERAND (strip_sections, 0); + { + /* The array-shaping operator only applies to a single pointer, not an + array section. */ + if (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION) + { + tree array = TREE_OPERAND (strip_sections, 0); + tree length = TREE_OPERAND (strip_sections, 2); + tree atype = TREE_TYPE (array); + tree eltype = atype != NULL_TREE ? TREE_TYPE (atype) : NULL_TREE; + if (eltype != NULL_TREE + && POINTER_TYPE_P (eltype) + && (length == NULL_TREE || !integer_onep (length))) + { + error ("OpenMP array shaping operator with non-pointer " + "argument"); + return error_mark_node; + } + } + strip_sections = TREE_OPERAND (strip_sections, 0); + } /* Determine the element type, either directly or by using "decltype" of an expression representing an element to figure it out later during template instantiation. */ if (type_dependent_expression_p (expr)) { + /* Peel ARRAY_TYPE layers before deferring the final dereference, + mirroring the non-dependent case below. Without this, a single + deferred dereference of the un-indexed base only strips one layer + via ordinary array-to-pointer decay, leaving any further array + dimensions in the resolved type, where they are later mistaken for an + omitted trailing shape dimension. (See also + gcc/testsuite/g++.dg/gomp/array-shaping-3.C). */ + tree btype = TREE_TYPE (strip_sections); + + tree base_expr = strip_sections; + if (btype != NULL_TREE) + { + if (TREE_CODE (btype) == REFERENCE_TYPE) + btype = TREE_TYPE (btype); + + while (TREE_CODE (btype) == ARRAY_TYPE) + { + base_expr + = build_min_nt_loc (loc, ARRAY_REF, base_expr, + integer_zero_node, NULL_TREE, NULL_TREE); + btype = TREE_TYPE (btype); + } + } + type = cxx_make_type (DECLTYPE_TYPE); DECLTYPE_TYPE_EXPR (type) - = build_min_nt_loc (loc, INDIRECT_REF, strip_sections); + = build_min_nt_loc (loc, INDIRECT_REF, base_expr); DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (type) = true; SET_TYPE_STRUCTURAL_EQUALITY (type); } @@ -13666,6 +13709,9 @@ cp_omp_create_arrayshape_type (location_t loc, tree expr, if (TREE_CODE (type) == REFERENCE_TYPE) type = TREE_TYPE (type); + while (TREE_CODE (type) == ARRAY_TYPE) + type = TREE_TYPE (type); + if (TREE_CODE (type) != POINTER_TYPE) { error ("OpenMP array shaping operator with non-pointer argument"); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index c0c3759e6da..2097cd3ffd1 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41635,9 +41635,10 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, location_t loc = token->location; decl = cp_parser_assignment_expression (parser); - if ((TREE_CODE (decl) == VIEW_CONVERT_EXPR - && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) - || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR) + if (TREE_TYPE (decl) != NULL_TREE + && ((TREE_CODE (decl) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR)) { reshaped_to = TREE_TYPE (decl); decl = TREE_OPERAND (decl, 0); @@ -41690,6 +41691,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, sections--; } + /* The previous loop may have uncovered a shaping cast. DECL_P + excludes such a cast over a further array section (reshaping + e.g. "x[1][1:2]"): that case is left for omp_expand_grid_dim's + SHAPE_TYPE handling in c-omp.cc. */ + if (!reshaped_to && TREE_TYPE (decl) != NULL_TREE + && TREE_CODE (decl) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE + && DECL_P (TREE_OPERAND (decl, 0))) + { + reshaped_to = TREE_TYPE (decl); + decl = TREE_OPERAND (decl, 0); + } + /* The handling of INDIRECT_REF here in the presence of array-shaping operations is a little tricky. We need to avoid treating a pointer dereference as a unit-sized array @@ -41764,7 +41778,54 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, } } - if (reshaped_to) + if (reshaped_to && TREE_TYPE (decl) != NULL_TREE + && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + { + unsigned array_layers = 0; + tree basetype = TREE_TYPE (decl); + while (TREE_CODE (basetype) == ARRAY_TYPE) + { + basetype = TREE_TYPE (basetype); + array_layers++; + } + + /* A further level of indirection here (e.g. "int **x[N]") + is not supported (for now); treating the second pointer's + bytes as the reshaped array's elements would silently read + the wrong data. (See also + gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C). */ + if (TREE_CODE (TREE_TYPE (basetype)) == POINTER_TYPE) + { + sorry_at (loc, "more than one level of pointer " + "indirection in a noncontiguous target " + "update"); + decl = error_mark_node; + } + else if (dims.length () != array_layers) + { + error_at (loc, "too many array section specifiers " + "for %qT", reshaped_to); + decl = error_mark_node; + } + else + { + for (unsigned i = 0; i < array_layers; i++) + { + omp_dim &d = dims[dims.length () - 1 - i]; + decl = grok_omp_array_section (loc, decl, + d.low_bound, d.length, + d.stride); + } + dims.truncate (0); + decl = cp_build_omp_arrayshape_cast (loc, reshaped_to, + decl, + tf_warning_or_error); + tree elems = array_type_nelts_total (reshaped_to); + decl = grok_omp_array_section (loc, decl, size_zero_node, + elems, NULL_TREE); + } + } + else if (reshaped_to) { unsigned reshaped_dims = 0; diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 4bf65d28fc9..16abf6b4b56 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -6053,15 +6053,16 @@ public: can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above case though, as some lengths could be zero. NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array - section. */ + section. TYPE carries the type of the inner section through recursion + (output only). */ static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, bool &non_contiguous, enum c_omp_region_type ort, - int *discontiguous) + int *discontiguous, tree *type) { - tree ret, low_bound, length, stride, type; + tree ret, low_bound, length, stride; bool openacc = (ort & C_ORT_ACC) != 0; if (TREE_CODE (t) != OMP_ARRAY_SECTION) { @@ -6113,6 +6114,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, if (type_dependent_expression_p (ret)) return NULL_TREE; ret = convert_from_reference (ret); + *type = TREE_TYPE (ret); return ret; } @@ -6124,14 +6126,10 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false); ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types, maybe_zero_len, first_non_one, - non_contiguous, ort, discontiguous); + non_contiguous, ort, discontiguous, type); if (ret == error_mark_node || ret == NULL_TREE) return ret; - if (TREE_CODE (ret) == OMP_ARRAY_SECTION) - type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0))); - else - type = TREE_TYPE (ret); low_bound = TREE_OPERAND (t, 1); length = TREE_OPERAND (t, 2); stride = TREE_OPERAND (t, 3); @@ -6235,11 +6233,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && (TREE_CODE (length) != INTEGER_CST || integer_onep (length))) first_non_one++; } - if (TREE_CODE (type) == ARRAY_TYPE) + if (TREE_CODE (*type) == ARRAY_TYPE) { if (length == NULL_TREE - && (TYPE_DOMAIN (type) == NULL_TREE - || TYPE_MAX_VALUE (TYPE_DOMAIN (type)) == NULL_TREE)) + && (TYPE_DOMAIN (*type) == NULL_TREE + || TYPE_MAX_VALUE (TYPE_DOMAIN (*type)) == NULL_TREE)) { error_at (OMP_CLAUSE_LOCATION (c), "for unknown bound array type length expression must " @@ -6263,13 +6261,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } - if (TYPE_DOMAIN (type) - && TYPE_MAX_VALUE (TYPE_DOMAIN (type)) - && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (type))) - == INTEGER_CST) + if (TYPE_DOMAIN (*type) && TYPE_MAX_VALUE (TYPE_DOMAIN (*type)) + && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (*type))) == INTEGER_CST) { tree size - = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (type))); + = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (*type))); size = size_binop (PLUS_EXPR, size, size_one_node); if (TREE_CODE (low_bound) == INTEGER_CST) { @@ -6296,11 +6292,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } maybe_zero_len = true; } - else if (length == NULL_TREE - && first_non_one == types.length () - && tree_int_cst_equal - (TYPE_MAX_VALUE (TYPE_DOMAIN (type)), - low_bound)) + else if (length == NULL_TREE && first_non_one == types.length () + && tree_int_cst_equal ( + TYPE_MAX_VALUE (TYPE_DOMAIN (*type)), low_bound)) first_non_one++; } else if (length == NULL_TREE) @@ -6382,7 +6376,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } } } - else if (TYPE_PTR_P (type)) + else if (TYPE_PTR_P (*type)) { if (length == NULL_TREE) { @@ -6459,7 +6453,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) - types.safe_push (type); + types.safe_push (*type); /* We will need to evaluate lb more than once. */ tree lb = cp_save_expr (low_bound); if (lb != low_bound) @@ -6488,6 +6482,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, else ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL, tf_warning_or_error); + + *type = TREE_TYPE (*type); return ret; } @@ -6528,9 +6524,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY) && OMP_ITERATOR_DECL_P (*tp)) tp = &TREE_VALUE (*tp); - tree first = handle_omp_array_sections_1 (c, *tp, types, - maybe_zero_len, first_non_one, - non_contiguous, ort, discontiguous); + tree type; + tree first + = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, + non_contiguous, ort, discontiguous, &type); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -6573,6 +6570,11 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, return false; bool higher_discontiguous = false; + /* Whether the first dimension uses the extended array section syntax, + i.e. whether it specifies a length or a stride. */ + bool first_extended = false; + /* Whether any dimension container except the last is a pointer. */ + bool non_last_pointer = false; for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; t = TREE_OPERAND (t, 0)) @@ -6638,10 +6640,15 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, full_span = true; } + tree op = TREE_OPERAND (t, 0); + bool shaped + = TREE_CODE (op) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (op)) == ARRAY_TYPE + && TREE_CODE (TREE_OPERAND (op, 0)) == OMP_ARRAY_SECTION; if (!integer_onep (stride) || (higher_discontiguous - && (!integer_zerop (low_bound) - || !full_span))) + && (!integer_zerop (low_bound) || !full_span)) + || shaped) *discontiguous = 2; if (!integer_onep (stride) @@ -6652,20 +6659,29 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, if (!maybe_zero_len && i > first_non_one) { + bool full_span; if (integer_nonzerop (low_bound)) goto is_noncontiguous; - if (length != NULL_TREE - && TREE_CODE (length) == INTEGER_CST - && TYPE_DOMAIN (types[i]) - && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])) - && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))) - == INTEGER_CST) + if (length != NULL_TREE && TREE_CODE (length) == INTEGER_CST) { - tree size; - size = size_binop (PLUS_EXPR, - TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])), - size_one_node); - if (!tree_int_cst_equal (length, size)) + /* A pointer-typed dimension has no static bound to compare + against, so it can never be proven to span the whole + dimension; treat it conservatively as not full span. */ + full_span = false; + if (TREE_CODE (types[i]) == ARRAY_TYPE + && TYPE_DOMAIN (types[i]) + && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])) + && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))) + == INTEGER_CST) + { + tree size; + size + = size_binop (PLUS_EXPR, + TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])), + size_one_node); + full_span = tree_int_cst_equal (length, size); + } + if (!full_span) { is_noncontiguous: if (discontiguous && *discontiguous) @@ -6743,9 +6759,26 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, else size = size_binop (MULT_EXPR, size, l); } + + if (i == 0 + && ((length != NULL_TREE && !integer_onep (length)) + || !integer_onep (stride))) + first_extended = true; + if (i < num - 1 && TREE_CODE (types[i]) == POINTER_TYPE + && /* accept decayed array */ !( + i == 0 && TREE_CODE (TREE_TYPE (types[i])) == ARRAY_TYPE)) + non_last_pointer = true; } + if (!processing_template_decl) { + if (first_extended && non_last_pointer) + error_at ( + OMP_CLAUSE_LOCATION (c), + "only the last dimension can refer to a pointer when the first " + "dimension uses the extended array section syntax in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + if (non_contiguous) { int kind = OMP_CLAUSE_MAP_KIND (c); @@ -10329,7 +10362,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_SHAPE_DIM) break; /* FALLTHRU */ case OMP_CLAUSE_TO: diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 5dd594ac211..ab723c1202c 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -4229,8 +4229,9 @@ omp_parse_noncontiguous_array (tree *expr0) tree expr = *expr0; bool noncontig = false; - while (TREE_CODE (expr) == OMP_ARRAY_SECTION - || TREE_CODE (expr) == ARRAY_REF) + while (TREE_CODE (expr) == OMP_ARRAY_SECTION || TREE_CODE (expr) == ARRAY_REF + || (TREE_CODE (expr) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_OPERAND (expr, 0)) == OMP_ARRAY_SECTION)) { /* Contiguous arrays use ARRAY_REF. By the time we reach here, OMP_ARRAY_SECTION is only used for noncontiguous arrays. */ diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-1.c b/gcc/testsuite/c-c++-common/gomp/array-section-1.c new file mode 100644 index 00000000000..99756bbf4c2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c @@ -0,0 +1,33 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* Test target update with a single level of pointer indirection, one dimension + before and after. */ + +#define DIM1 5 +#define DIM2 10 + +void fixed_index (void) +{ + int *x[DIM1]; +#pragma omp target update to(x[2][ :DIM2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */ +} + +/* Index, length and stride need not be literal constants -- a variable + works just as well. */ + +void range_of_pointers (int lo, int len) +{ + int *x[DIM1]; +#pragma omp target update to(x[lo:len][ :DIM2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:SAVE_EXPR \[len: len\]\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */ +} + +void strided_range_of_pointers (int lo, int len, int str) +{ + int *x[DIM1]; +#pragma omp target update to(x[lo:len:str][ :DIM2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:SAVE_EXPR \[len: len\]\) map\(grid_stride:str\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-2.c b/gcc/testsuite/c-c++-common/gomp/array-section-2.c new file mode 100644 index 00000000000..15bc2cb28ba --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-2.c @@ -0,0 +1,15 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* Test target update with a 2D array of pointers */ + +#define DIM1 5 +#define DIM2 5 + +void two_d_array_of_pointers (void) +{ + int *y[DIM1][DIM2]; +#pragma omp target update to(y[0][ :2]) +/* { dg-final { scan-tree-dump {map\(to_grid:y \[len: [0-9]+\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:0 \[len: 2\]\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-3.c b/gcc/testsuite/c-c++-common/gomp/array-section-3.c new file mode 100644 index 00000000000..f7b17475b5d --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-3.c @@ -0,0 +1,38 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* A 2D array-shaping cast applied to a 2D array of pointers -- two segments: + the array-of-pointers side, selected by two fixed indices, and the reshaped + pointee side, both of its dimensions explicitly sectioned. */ + +#define DIM1 3 +#define DIM2 4 +#define ROWS 6 +#define COLS 6 + +void shape_cast_two_fixed_indices_strided_row (void) +{ + int *x[DIM1][DIM2]; +#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2][0:2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } */ +} + +void shape_cast_length_one_range_index (void) +{ + int *x[DIM1][DIM2]; +#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2][1:3]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 2\] \[pointer\]\) map\(shape_dim:6\) map\(grid_dim:1 \[len: 3\]\) map\(shape_dim:6\)} "original" } } */ +} + +/* Same shape as case 1, but with variables everywhere a literal is + allowed: the two fixed array-of-pointers indices, and the reshaped + pointee side's index/length/stride. */ + +void shape_cast_var (int i, int j, int r0, int rlen, int rstr, int c0, + int clen) +{ + int *x[DIM1][DIM2]; +#pragma omp target update to((([ROWS][COLS]) x[i][j])[r0:rlen:rstr][c0:clen]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:i \[len: 1\]\) map\(grid_dim:j \[len: 1\]\) map\(grid_dim:SAVE_EXPR \[len: rlen\] \[pointer\]\) map\(grid_stride:rstr\) map\(shape_dim:6\) map\(grid_dim:SAVE_EXPR \[len: clen\]\) map\(shape_dim:6\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-4.c b/gcc/testsuite/c-c++-common/gomp/array-section-4.c new file mode 100644 index 00000000000..ce446b931c2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-4.c @@ -0,0 +1,21 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* An array-of-pointers with several dimensions selected by consecutive fixed + indices before the final range. + Confirms walking several plain ARRAY_REFs (rather than OMP_ARRAY_SECTIONs) + before reaching the pointer dereference produces one GOMP_MAP_GRID_DIM per + fixed index, plus a final one for the pointer crossing. */ + +#define D1 3 +#define D2 3 +#define D3 4 +#define M 8 + +void three_consecutive_fixed_indices (void) +{ + int *w[D1][D2][D3]; +#pragma omp target update to(w[1][2][0][2:4]) +/* { dg-final { scan-tree-dump {map\(to_grid:w \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 4\] \[pointer\]\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-5.c b/gcc/testsuite/c-c++-common/gomp/array-section-5.c new file mode 100644 index 00000000000..9bd8a251393 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-5.c @@ -0,0 +1,25 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* Test an array-shaping cast that has more dimensions than are explicitly + sectioned afterwards -- copy the whole span gathered from the cast. */ + +#define DIM1 3 +#define DIM2 4 +#define ROWS 6 +#define COLS 6 + +void shape_cast_omitted_column_dim_strided (void) +{ + int *x[DIM1][DIM2]; +#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 6\]\) map\(shape_dim:6\)} "original" } } */ +} + +void shape_cast_omitted_column_dim_length_one_range (void) +{ + int *x[DIM1][DIM2]; +#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 2\] \[pointer\]\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 6\]\) map\(shape_dim:6\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-6.c b/gcc/testsuite/c-c++-common/gomp/array-section-6.c new file mode 100644 index 00000000000..e90f4e66188 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-6.c @@ -0,0 +1,18 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* An array-shaping cast applied to a plain pointer (no array-of-pointers + indirection involved), with a further, unit-stride, contiguous section + applied outside the cast. Check that it simplifies correctly and does not + go through the noncontiguous-update codepath. */ + +#define N 100 + +void plain_pointer_shape_cast (void) +{ + int *ptr; +#pragma omp target update to((([N]) ptr)[10:30]) +/* { dg-final { scan-tree-dump {to\(VIEW_CONVERT_EXPR\(\*ptr\)\[10\] \[len: 120\]\)} "original" } } */ +/* { dg-final { scan-tree-dump-not "to_grid" "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-8.c b/gcc/testsuite/c-c++-common/gomp/array-section-8.c new file mode 100644 index 00000000000..4873138362a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-8.c @@ -0,0 +1,27 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* Check three-segment (two pointer indirections) noncontiguous update: a + strided section reached by crossing two levels of pointer indirection. */ + +#define DIM1 4 + +void three_segment_one_dim_per_segment (void) +{ + int **arr[DIM1]; +#pragma omp target update from(arr[2][0:2:2][3:5:2]) +/* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 5\] \[pointer\]\) map\(grid_stride:2\)} "original" } } */ +} + +/* Same shape, but every index/length/stride across all three segments is + a variable rather than a literal. */ + +void three_segment_one_dim_per_segment_var (int p, int p0, int plen, + int pstr, int e0, int elen, + int estr) +{ + int **arr[DIM1]; +#pragma omp target update from(arr[p][p0:plen:pstr][e0:elen:estr]) +/* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:SAVE_EXPR

\[len: 1\]\) map\(grid_dim:SAVE_EXPR \[len: plen\] \[pointer\]\) map\(grid_stride:pstr\) map\(grid_dim:SAVE_EXPR \[len: elen\] \[pointer\]\) map\(grid_stride:estr\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-9.c b/gcc/testsuite/c-c++-common/gomp/array-section-9.c new file mode 100644 index 00000000000..9d2e0ceae09 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/array-section-9.c @@ -0,0 +1,21 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */ + +/* Check three-segment (two pointer indirections) noncontiguous update with two + dimensions selected per segment -- six dimensions total. */ + +#define DIM1A 2 +#define DIM1B 3 +#define ROWLEN 20 +#define NCOLS 3 + +typedef int (*p2_t)[ROWLEN]; +typedef p2_t (*p1_t)[NCOLS]; + +void three_segment_two_dims_per_segment (void) +{ + p1_t arr[DIM1A][DIM1B]; +#pragma omp target update from(arr[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2]) +/* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:0 \[len: 2\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_dim:3 \[len: 5\]\) map\(grid_stride:2\)} "original" } } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c new file mode 100644 index 00000000000..fa2fccd93c6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c @@ -0,0 +1,19 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +/* Regression test for strided target updates with iterators. */ + +#define DIM1 17 +#define DIM2 39 + +void f (int *x[DIM1]) +{ + /* Per-index iterator selecting the pointer, no stride on the pointee: + still supported. */ +#pragma omp target update to (iterator(i=0:DIM1): x[i][ :DIM2]) + + /* Per-index iterator selecting the pointer, but an explicit stride on + the pointee: still sorry. */ +#pragma omp target update to (iterator(i=0:DIM1): x[i][0:DIM2:2]) + /* { dg-message "sorry, unimplemented: strided target updates with iterators" "" { target *-*-* } .-1 } */ +} diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-3.C b/gcc/testsuite/g++.dg/gomp/array-shaping-3.C new file mode 100644 index 00000000000..423b298d4ff --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/array-shaping-3.C @@ -0,0 +1,71 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +/* Test target update on array-shaping cast applied to a section selected from a + 2D array-of-pointers base, when the enclosing function is a template. */ + +#define DIM1 7 +#define DIM2 9 +#define ROWS 6 +#define COLS 6 + +template +void foo() +{ + T *x[DIM1][DIM2]; + +#pragma omp target update to((([ROWS][COLS]) x[1][2])[1:3:2][0:2]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } +// { dg-final { scan-tree-dump-not {grid_dim:0 \[len: 9\]} "original" } } +// { dg-final { scan-tree-dump-not {shape_dim:9} "original" } } +} + +template +void bar() +{ + T *x[DIM1]; + +#pragma omp target update to((([ROWS][COLS]) x[1])[1:3:2][0:2]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } +// { dg-final { scan-tree-dump-not {grid_dim:0 \[len: 9\]} "original" } } +// { dg-final { scan-tree-dump-not {shape_dim:9} "original" } } +} + +void control() +{ + int *x[DIM1][DIM2]; + +#pragma omp target update to((([ROWS][COLS]) x[1][2])[1:3:2][0:2]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } +} + +/* Reference + template combination: same 2D array-of-pointers + base as foo() above, but accessed through a C++ reference. */ + +template +void baz() +{ + T *x[DIM1][DIM2]; + T *(&x_ref)[DIM1][DIM2] = x; + +#pragma omp target update to((([ROWS][COLS]) x_ref[1][2])[1:3:2][0:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*x_ref \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } +// { dg-final { scan-tree-dump-not {grid_dim:0 \[len: 9\]} "original" } } +// { dg-final { scan-tree-dump-not {shape_dim:9} "original" } } +} + +void control_ref() +{ + int *x[DIM1][DIM2]; + int *(&x_ref)[DIM1][DIM2] = x; + +#pragma omp target update to((([ROWS][COLS]) x_ref[1][2])[1:3:2][0:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*x_ref \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } +} + +void g() +{ + foo (); + bar (); + baz (); +} diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-4.C b/gcc/testsuite/g++.dg/gomp/array-shaping-4.C new file mode 100644 index 00000000000..e857ff01c1e --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/array-shaping-4.C @@ -0,0 +1,37 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +/* Test an array-shaping cast applied directly to a single element of an + array of pointers, with and without outer array section. + Also check templates and variable subscript indices. */ + +#define DIM1 5 +#define N 20 + +template +void foo() +{ + T *x[DIM1]; + +#pragma omp target update to(([N]) x[2]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 20\] \[pointer\]\) map\(shape_dim:20\)} "original" } } + +#pragma omp target update to((([N][N+1]) x[2])[1:3:2][1:4]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:20\) map\(grid_dim:1 \[len: 4\]\) map\(shape_dim:21\)} "original" } } +} + +void control() +{ + int *x[DIM1]; + +#pragma omp target update to(([N]) x[2]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 20\] \[pointer\]\) map\(shape_dim:20\)} "original" } } + +#pragma omp target update to((([N][N+1]) x[2])[1:3:2][1:4]) +// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:20\) map\(grid_dim:1 \[len: 4\]\) map\(shape_dim:21\)} "original" } } +} + +void g() +{ + foo (); +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C new file mode 100644 index 00000000000..c8ce6f61fe0 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C @@ -0,0 +1,51 @@ +// { dg-do compile } + +/* Test rejected array-shaping cast combinations. */ + +#define DIM1 5 +#define DIM2 5 + +template +void foo () +{ + T *x[DIM1]; + +#pragma omp target update to(([DIM2]) x[1:3][0:5]) +// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } +// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + + T **y[DIM1]; + +#pragma omp target update to(([DIM2]) y[0:2]) +// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } +// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + + /* A *fixed* index into the same array-of-pointers-of-pointers -- unlike the range above, this is not rejected by the non-pointer-argument check: a fixed index does denote a single pointer expression. But reshaping it here would still leave a pointer-typed result ("T *[DIM2]"): copying that host-to-device as plain data would copy raw host pointer values with no attach translation, which is unsupported (for now). */ +#pragma omp target update to(([DIM2]) y[2]) + // { dg-message "sorry, unimplemented: more than one level of pointer indirection in a noncontiguous target update" "" { target *-*-* } .-1 } + // { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } +} + +void control () +{ + int *x[DIM1]; + +#pragma omp target update to(([DIM2]) x[1:3][0:5]) +// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } +// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + + int **y[DIM1]; + +#pragma omp target update to(([DIM2]) y[0:2]) +// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } +// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + +#pragma omp target update to(([DIM2]) y[2]) +// { dg-message "sorry, unimplemented: more than one level of pointer indirection in a noncontiguous target update" "" { target *-*-* } .-1 } +// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } +} + +void g () +{ + foo (); +} diff --git a/gcc/testsuite/gcc.dg/gomp/array-shaping-9.c b/gcc/testsuite/gcc.dg/gomp/array-shaping-9.c new file mode 100644 index 00000000000..9050a56515a --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/array-shaping-9.c @@ -0,0 +1,25 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + + +#define DIM1 5 +#define N 20 + +int main () +{ + int *x[DIM1]; + + /* An array-shaping cast applied directly to a single element of an + array of pointers, with no further outer section. */ + +#pragma omp target update to(([N]) x[2]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 20\] \[pointer\]\) map\(shape_dim:20\)} "original" } } */ + + /* Same base, but with an explicit outer section applied to the + reshaped 2D result. */ + +#pragma omp target update to((([N][N+1]) x[2])[1:3:2][1:4]) +/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:20\) map\(grid_dim:1 \[len: 4\]\) map\(shape_dim:21\)} "original" } } */ + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c b/gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c new file mode 100644 index 00000000000..e2bf421baf8 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c @@ -0,0 +1,60 @@ +/* { dg-do compile } */ + +/* Test for bad array section syntax in OMP target update clauses implying + pointer indirection. */ + +#include + +#define DIM1 5 +#define DIM2 5 + +int main (void) +{ + int **z; + + z = malloc(DIM1 * sizeof z); + for (int i = 0; i < DIM1; i++) + z[i] = malloc(DIM2 * sizeof(int)); + + #pragma omp target update to(z[:DIM1][0]) + /* { dg-error "only the last dimension can refer to a pointer when the first dimension uses the extended array section syntax" "" { target *-*-* } .-1 } */ + + return 0; +} + +#define D1 3 +#define D2 4 +#define ROWS 6 +#define COLS 6 + +void +g (void) +{ + int *x[D1][D2]; + int i, j; + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + x[i][j] = malloc (ROWS * COLS * sizeof (int)); + + #pragma omp target update to(x[1][1:2][1:3][0:2]) + /* { dg-error "does not have pointer or array type" "" { target *-*-* } .-1 } */ + /* { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */ +} + +void +h (void) +{ + int **y[D1]; + int i, j; + + for (i = 0; i < D1; i++) + { + y[i] = malloc (D2 * sizeof (int *)); + for (j = 0; j < D2; j++) + y[i][j] = malloc (10 * sizeof (int)); + } + + #pragma omp target update to(y[0:2][0:3][0:5]) + /* { dg-error "only the last dimension can refer to a pointer when the first dimension uses the extended array section syntax" "" { target *-*-* } .-1 } */ +} diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c new file mode 100644 index 00000000000..b05b182f2f5 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c @@ -0,0 +1,45 @@ +/* { dg-do compile } */ + +/* Test rejected array-shaping cast combinations. */ + +#define DIM1 5 +#define DIM2 5 + +int main (void) +{ + int *x[DIM1]; + + /* Array-shaping operator applied on an array section. */ +#pragma omp target update to(([DIM2]) x[1:3][0:5]) + /* { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } */ + + return 0; +} + +/* Same restriction, on an array-of-pointers-of-pointers. */ + +void +depth2 (void) +{ + int **y[DIM1]; + +#pragma omp target update to(([DIM2]) y[0:2]) +/* { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } */ +} + +/* A *fixed* index into the same array-of-pointers-of-pointers, e.g. + "y[2]" -- unlike the range above, this is not rejected by the + non-pointer-argument check. But reshaping it here would still leave + a pointer-typed result ("int *[DIM2]"): copying that host-to-device + as plain data would copy raw host pointer values with no attach + translation. */ + +void +depth2_fixed_index (void) +{ + int **y[DIM1]; + +#pragma omp target update to(([DIM2]) y[2]) + /* { dg-message "sorry, unimplemented: more than one level of pointer indirection in a noncontiguous target update" "" { target *-*-* } .-1 } */ + /* { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */ +} From patchwork Wed Aug 5 20:23:41 2026 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 140678 Return-Path: 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 2AFF14BAE7E6 for ; Wed, 5 Aug 2026 20:26:37 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 2AFF14BAE7E6 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=h2KVD9mV X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wr1-x432.google.com (mail-wr1-x432.google.com [IPv6:2a00:1450:4864:20::432]) by sourceware.org (Postfix) with ESMTPS id 791E94BAE7E8 for ; Wed, 5 Aug 2026 20:25:42 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 791E94BAE7E8 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 791E94BAE7E8 Authentication-Results: sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::432 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961542; cv=none; b=i0VVH9zOeU4ZM/P6cmFPMEOeU4E1T9iY0sD2qElV3BS2HBZWY/TX32FsWUNEUKe0kQDdut3JoNttyUdp2S7v8wm1QdCi3cKsVM+x+8E2LbFGP0sD0Il3T4omE91wNG2jGtsgW1ode3OootzqXD54R6c6S8R54FjeBSgcf40bl0I= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961542; c=relaxed/simple; bh=TbTtL1jc6ykNJSAmu7Sd/rLy5OcOLrVmYyl8th4KoeU=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=tO6We/BIbKY4hNRpBqKfsDgWOcBpDOgDsyugP9dn+wg3KjHQp6te9Q3RA+kBqpyr/SU/HMIJ/WV1JNsfwG2IpT32GRISidxED4lC/133XKzRAKIqwfYzrZl5OTPJ3Lg5vIWn0OahGJvcyBsanOX8/fXLzSAlMaEHRKm67wOpVIo= 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=h2KVD9mV DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 791E94BAE7E8 Received: by mail-wr1-x432.google.com with SMTP id ffacd0b85a97d-476a130c138so1310885f8f.0 for ; Wed, 05 Aug 2026 13:25:42 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1785961541; x=1786566341; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to:content-type; bh=2L9hdx/6XOJpG1ivZykNI8seUTNK2CCG7yxUCVayOwI=; b=h2KVD9mVZvLfqsSgqm576bqZQ8v9tOamVmGlmMfpccRk/89wBaR1BZsuWqs36bb6rw j+M+fE+xQ06clRzU4A1+boVM91l1SyaMqMUS6wCSPVPDYgjPw86aXOm4+2hHUEi5kW/g pPx5YoQk6Xh+/IKb7/8FhVCHtBQnHxw/PK2qfVh7IbjUahQNaLTEwfxQGD0q1//46oD6 2tU27AQwzHqIBqJQWSnA+6yVYBAJML8BEqRCWC+iMa093Se4OADxlCrSJVshGsNXv6If mzLOFf0oFsI5XrORAKtLakdk+bE/YxsbeLjHaWNFnauLfQcyRKiW685D6T1hhTrJNonZ tOvw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1785961541; x=1786566341; h=content-transfer-encoding:mime-version:references:in-reply-to :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=2L9hdx/6XOJpG1ivZykNI8seUTNK2CCG7yxUCVayOwI=; b=Qjk77R9PGKdAjZ99ScXxMtOvZ9v9K7WL4AUyLaIQuY/t01fkjzBMDXJ7y0JgITn7md iCst/jRwqoc6VLMkUeWVAxZqDw3vgoMz3TA713bIiW8T0jIId23T5Bf5iLn39bpJZlGO BYKYLNEnBYhhillJC7SYYcsOjGTJUpidYNEIHRS3yUDQ4ghREd8SgKv3Znhv+kQsIxW+ onnrJA7BllgqU9AAINq94Vle4SFJNirwd7W/fCJQYxtzWKaavW7oDZk6CP/HCmtpWaFb sNUXhhVTnOvikhaU1frIOe+411fF3XhaVkCWMljltjqSA+g0LC4E0tt+BMcSPrdPwB99 z2jg== X-Gm-Message-State: AOJu0YzRDo6Q/7sww+2TWsepJaml6hMn6OeOkpPlYpU2CqD4qYyu86OA 4e3iDih6m7+g/H1g/IG2AuRo2KaHxDOEA5jOudRPs43Ee/0c/LnTmcXNCsCjr/POyfEu6hNe17U hCgtz X-Gm-Gg: AR+sD11NtAURDh1LR+/lN2SIrTlfpB8zUbenC5Ik8mQDZmf25fFRlnNzL3obIRVS1kU nrXCdefEInFzM8U5TVtV+NeYNuXhNA/TDVFzRTHraxff3sfb70QRnZVSPzO7riGMUx44jQTXsbr 2CgZRqHKND5uldn6aYpZXV3xj2By7horqxbhfMjLJqDA3wWgL7Z1fj80jLvxBJgof+0On/CUjlC H7gAYYt5GAtOuNNu8raQnrcqVVGAOk7IGy6upC3n4O8tUrsVDDqz9xwKq27hS17/YqfKoVWcqJo xWQb+lw+oA273XVw7N19AB93hqYn714pjEbKp4Yll5TPxf/2YzMnb7rk4lNMAJENneg5KHNqawQ nCHqk3y/oWfuou1AXKey9L3+X+mfMWNyVsYnCRG27VQPZ2Zy2XC6jBsaGNDHlVWck6wo9KV0dbv MI7wT/wk8EqZcDhNrW6vJ3WgXt3Hg9sKatvZ4NA6GDJomWZe2/vGCfgw== X-Received: by 2002:a05:6000:46cb:b0:46e:1815:6a83 with SMTP id ffacd0b85a97d-47fec62c565mr11998268f8f.29.1785961541189; Wed, 05 Aug 2026 13:25:41 -0700 (PDT) Received: from raptor ([62.108.198.223]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-47ff79b431dsm67029f8f.13.2026.08.05.13.25.39 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 05 Aug 2026 13:25:40 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, Paul-Antoine Arras Subject: [PATCH 3/5] openmp: Outline lower_omp_target_grid_desc from lower_omp_target Date: Wed, 5 Aug 2026 22:23:41 +0200 Message-ID: <20260805202343.2868178-4-parras@baylibre.com> X-Mailer: git-send-email 2.53.0 In-Reply-To: <20260805202343.2868178-1-parras@baylibre.com> References: <20260805202343.2868178-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-12.8 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, 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 List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org Move the noncontiguous-array descriptor construction for GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clauses out of lower_omp_target's clause-processing loop into its own static function. Pure code motion, no behavior change; the generalization to array-shaping casts, array-of-pointers sections, and multi-segment descriptors follows in a separate change. gcc/ChangeLog: * omp-low.cc (lower_omp_target_grid_desc): New, extracted from lower_omp_target. (lower_omp_target): Call it instead of building the descriptor inline. --- gcc/omp-low.cc | 656 +++++++++++++++++++++++++------------------------ 1 file changed, 335 insertions(+), 321 deletions(-) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 837a0493c91..45f6c701f48 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -13723,6 +13723,340 @@ convert_from_firstprivate_int (tree var, tree orig_type, bool is_ref, gimplify_assign (tmp, var, gs); return fold_build1 (VIEW_CONVERT_EXPR, type, tmp); + +/* Build the noncontiguous-array descriptor for the + GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause C (whose GOMP_MAP_TO_PSET + sibling clause holds the descriptor decl and the GOMP_MAP_GRID_DIM/ + GOMP_MAP_GRID_STRIDE clauses that follow describe each dimension). Any + gimplification needed to initialize the descriptor is appended to + *ILIST_P. */ + +static void +lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) +{ + tree decl = OMP_CLAUSE_DECL (c); + tree dn = OMP_CLAUSE_CHAIN (c); + gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET); + tree desc = OMP_CLAUSE_DECL (dn); + + tree oc, elsize = OMP_CLAUSE_SIZE (c); + tree type = TREE_TYPE (decl); + int i, dims = 0; + tree nc; + auto_vec tdims; + bool pointer_based = false, handled_pointer_section = false; + tree arrsize = size_one_node; + + /* Allow a single (maybe strided) array section if we have a + pointer base. */ + if (TREE_CODE (decl) == INDIRECT_REF + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == POINTER_TYPE)) + { + pointer_based = true; + dims = 1; + } + else + /* NOTE: Don't treat (e.g. Fortran, fixed-length) strings as + array types here; array section syntax isn't applicable to + strings. */ + for (tree itype = type; + TREE_CODE (itype) == ARRAY_TYPE + && !TYPE_STRING_FLAG (itype); + itype = TREE_TYPE (itype)) + { + tdims.safe_push (itype); + dims++; + } + + unsigned tdim = 0; + + vec *vdim; + vec *vindex; + vec *vlen; + vec *vstride; + vec_alloc (vdim, dims); + vec_alloc (vindex, dims); + vec_alloc (vlen, dims); + vec_alloc (vstride, dims); + + tree size_arr_type + = build_array_type_nelts (size_type_node, dims); + + tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim"); + DECL_NAMELESS (dim_tmp) = 1; + TREE_ADDRESSABLE (dim_tmp) = 1; + TREE_STATIC (dim_tmp) = 1; + tree index_tmp = create_tmp_var (size_arr_type, ".omp_index"); + DECL_NAMELESS (index_tmp) = 1; + TREE_ADDRESSABLE (index_tmp) = 1; + TREE_STATIC (index_tmp) = 1; + tree len_tmp = create_tmp_var (size_arr_type, ".omp_len"); + DECL_NAMELESS (len_tmp) = 1; + TREE_ADDRESSABLE (len_tmp) = 1; + TREE_STATIC (len_tmp) = 1; + tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride"); + DECL_NAMELESS (stride_tmp) = 1; + TREE_ADDRESSABLE (stride_tmp) = 1; + TREE_STATIC (stride_tmp) = 1; + + oc = c; + c = dn; + + tree span = NULL_TREE; + + for (i = 0; i < dims; i++) + { + nc = OMP_CLAUSE_CHAIN (c); + tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE, + stride = size_one_node; + + if (nc + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM) + { + index = OMP_CLAUSE_DECL (nc); + len = OMP_CLAUSE_SIZE (nc); + + index = fold_convert (sizetype, index); + len = fold_convert (sizetype, len); + + tree nc2 = OMP_CLAUSE_CHAIN (nc); + if (nc2 + && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc2) + == GOMP_MAP_GRID_STRIDE)) + { + stride = OMP_CLAUSE_DECL (nc2); + stride = fold_convert (sizetype, stride); + if (OMP_CLAUSE_SIZE (nc2)) + { + /* If the element size is not the same as the + distance between two adjacent array + elements (in the innermost dimension), + retrieve the latter value ("span") from the + size field of the stride. We only expect to + see one such field per array. */ + gcc_assert (!span); + span = OMP_CLAUSE_SIZE (nc2); + span = fold_convert (sizetype, span); + } + nc = nc2; + } + + 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 + in libgomp. */ + tree dtype = TYPE_DOMAIN (tdims[tdim]); + tree minval = TYPE_MIN_VALUE (dtype); + tree maxval = TYPE_MAX_VALUE (dtype); + minval = fold_convert (sizetype, minval); + maxval = fold_convert (sizetype, maxval); + dim = size_binop (MINUS_EXPR, maxval, minval); + dim = size_binop (PLUS_EXPR, dim, + size_one_node); + arrsize = size_binop (MULT_EXPR, arrsize, dim); + } + else if (pointer_based && !handled_pointer_section) + { + /* Use the selected array section to determine the + size of the array. */ + tree tmp = size_binop (MULT_EXPR, len, stride); + tmp = size_binop (MINUS_EXPR, tmp, stride); + tmp = size_binop (PLUS_EXPR, tmp, size_one_node); + dim = size_binop (PLUS_EXPR, index, tmp); + arrsize = size_binop (MULT_EXPR, arrsize, dim); + handled_pointer_section = true; + } + else + { + if (pointer_based) + error_at (OMP_CLAUSE_LOCATION (c), + "too many array section specifiers " + "for pointer-based array"); + else + error_at (OMP_CLAUSE_LOCATION (c), + "too many array section specifiers " + "for array"); + dim = index = len = stride = error_mark_node; + } + tdim++; + + c = nc; + } + else + { + /* We have more array dimensions than array section + specifiers. Copy the whole span. */ + tree dtype = TYPE_DOMAIN (tdims[tdim]); + tree minval = TYPE_MIN_VALUE (dtype); + tree maxval = TYPE_MAX_VALUE (dtype); + minval = fold_convert (sizetype, minval); + maxval = fold_convert (sizetype, maxval); + dim = size_binop (MINUS_EXPR, maxval, minval); + dim = size_binop (PLUS_EXPR, dim, size_one_node); + len = dim; + index = minval; + nc = c; + } + + if (TREE_CODE (dim) != INTEGER_CST) + TREE_STATIC (dim_tmp) = 0; + + if (TREE_CODE (index) != INTEGER_CST) + TREE_STATIC (index_tmp) = 0; + + if (TREE_CODE (len) != INTEGER_CST) + TREE_STATIC (len_tmp) = 0; + + if (TREE_CODE (stride) != INTEGER_CST) + TREE_STATIC (stride_tmp) = 0; + + tree cidx = size_int (i); + CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim); + CONSTRUCTOR_APPEND_ELT (vindex, cidx, index); + CONSTRUCTOR_APPEND_ELT (vlen, cidx, len); + CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride); + } + + tree bias = size_zero_node; + tree volume = size_one_node; + tree enclosure = size_one_node; + for (i = dims - 1; i >= 0; i--) + { + tree dim = (*vdim)[i].value; + tree index = (*vindex)[i].value; + tree stride = (*vstride)[i].value; + tree len = (*vlen)[i].value; + + /* For the bias we want, e.g.: + + index[0] * stride[0] * dim[1] * dim[2] + + index[1] * stride[1] * dim[2] + + index[2] * stride[2] + + All multiplied by "span" (or "elsize"). */ + + tree index_stride = size_binop (MULT_EXPR, index, stride); + bias = size_binop (PLUS_EXPR, bias, + size_binop (MULT_EXPR, volume, + index_stride)); + volume = size_binop (MULT_EXPR, volume, dim); + + if (i == 0) + { + tree elems_covered = size_binop (MINUS_EXPR, len, + size_one_node); + elems_covered = size_binop (MULT_EXPR, elems_covered, + stride); + elems_covered = size_binop (PLUS_EXPR, elems_covered, + size_one_node); + enclosure = size_binop (MULT_EXPR, enclosure, + elems_covered); + } + else + enclosure = volume; + } + + /* If we don't have a separate span size, use the element size + instead. */ + if (!span) + span = fold_convert (sizetype, elsize); + + /* The size of a volume enclosing the elements to be + transferred. */ + OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, enclosure, span); + /* And the bias of the first element we will update. */ + OMP_CLAUSE_SIZE (dn) = size_binop (MULT_EXPR, bias, span); + + tree cdim = build_constructor (size_arr_type, vdim); + tree cindex = build_constructor (size_arr_type, vindex); + tree clen = build_constructor (size_arr_type, vlen); + tree cstride = build_constructor (size_arr_type, vstride); + + if (TREE_STATIC (dim_tmp)) + DECL_INITIAL (dim_tmp) = cdim; + else + gimplify_assign (dim_tmp, cdim, ilist_p); + + if (TREE_STATIC (index_tmp)) + DECL_INITIAL (index_tmp) = cindex; + else + gimplify_assign (index_tmp, cindex, ilist_p); + + if (TREE_STATIC (len_tmp)) + DECL_INITIAL (len_tmp) = clen; + else + gimplify_assign (len_tmp, clen, ilist_p); + + if (TREE_STATIC (stride_tmp)) + DECL_INITIAL (stride_tmp) = cstride; + else + gimplify_assign (stride_tmp, cstride, ilist_p); + + tree desc_type = TREE_TYPE (desc); + + tree ndims_field = TYPE_FIELDS (desc_type); + tree elemsize_field = DECL_CHAIN (ndims_field); + tree span_field = DECL_CHAIN (elemsize_field); + tree dim_field = DECL_CHAIN (span_field); + tree index_field = DECL_CHAIN (dim_field); + tree len_field = DECL_CHAIN (index_field); + tree stride_field = DECL_CHAIN (len_field); + + vec *v; + vec_alloc (v, 7); + + bool all_static = (TREE_STATIC (dim_tmp) + && TREE_STATIC (index_tmp) + && TREE_STATIC (len_tmp) + && TREE_STATIC (stride_tmp)); + + dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node, + NULL_TREE, NULL_TREE); + dim_tmp = build_fold_addr_expr (dim_tmp); + + /* TODO: we could skip all-zeros index. */ + index_tmp = build4 (ARRAY_REF, sizetype, index_tmp, + size_zero_node, NULL_TREE, NULL_TREE); + index_tmp = build_fold_addr_expr (index_tmp); + + len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node, + NULL_TREE, NULL_TREE); + len_tmp = build_fold_addr_expr (len_tmp); + + /* TODO: we could skip all-ones stride. */ + stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp, + size_zero_node, NULL_TREE, NULL_TREE); + stride_tmp = build_fold_addr_expr (stride_tmp); + + elsize = fold_convert (sizetype, elsize); + tree ndims = size_int (dims); + + CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims); + CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize); + CONSTRUCTOR_APPEND_ELT (v, span_field, span); + CONSTRUCTOR_APPEND_ELT (v, dim_field, dim_tmp); + CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp); + CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp); + CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp); + + tree desc_ctor = build_constructor (desc_type, v); + + if (all_static) + { + TREE_STATIC (desc) = 1; + DECL_INITIAL (desc) = desc_ctor; + } + else + gimplify_assign (desc, desc_ctor, ilist_p); + + OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc); +} + } /* Lower the GIMPLE_OMP_TARGET in the current statement @@ -14427,327 +14761,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID)) { - tree decl = OMP_CLAUSE_DECL (c); - tree dn = OMP_CLAUSE_CHAIN (c); - gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET); - tree desc = OMP_CLAUSE_DECL (dn); - - tree oc, elsize = OMP_CLAUSE_SIZE (c); - tree type = TREE_TYPE (decl); - int i, dims = 0; - auto_vec tdims; - bool pointer_based = false, handled_pointer_section = false; - tree arrsize = size_one_node; - - /* Allow a single (maybe strided) array section if we have a - pointer base. */ - if (TREE_CODE (decl) == INDIRECT_REF - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == POINTER_TYPE)) - { - pointer_based = true; - dims = 1; - } - else - /* NOTE: Don't treat (e.g. Fortran, fixed-length) strings as - array types here; array section syntax isn't applicable to - strings. */ - for (tree itype = type; - TREE_CODE (itype) == ARRAY_TYPE - && !TYPE_STRING_FLAG (itype); - itype = TREE_TYPE (itype)) - { - tdims.safe_push (itype); - dims++; - } - - unsigned tdim = 0; - - vec *vdim; - vec *vindex; - vec *vlen; - vec *vstride; - vec_alloc (vdim, dims); - vec_alloc (vindex, dims); - vec_alloc (vlen, dims); - vec_alloc (vstride, dims); - - tree size_arr_type - = build_array_type_nelts (size_type_node, dims); - - tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim"); - DECL_NAMELESS (dim_tmp) = 1; - TREE_ADDRESSABLE (dim_tmp) = 1; - TREE_STATIC (dim_tmp) = 1; - tree index_tmp = create_tmp_var (size_arr_type, ".omp_index"); - DECL_NAMELESS (index_tmp) = 1; - TREE_ADDRESSABLE (index_tmp) = 1; - TREE_STATIC (index_tmp) = 1; - tree len_tmp = create_tmp_var (size_arr_type, ".omp_len"); - DECL_NAMELESS (len_tmp) = 1; - TREE_ADDRESSABLE (len_tmp) = 1; - TREE_STATIC (len_tmp) = 1; - tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride"); - DECL_NAMELESS (stride_tmp) = 1; - TREE_ADDRESSABLE (stride_tmp) = 1; - TREE_STATIC (stride_tmp) = 1; - - oc = c; - c = dn; - - tree span = NULL_TREE; - - for (i = 0; i < dims; i++) - { - nc = OMP_CLAUSE_CHAIN (c); - tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE, - stride = size_one_node; - - if (nc - && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM) - { - index = OMP_CLAUSE_DECL (nc); - len = OMP_CLAUSE_SIZE (nc); - - index = fold_convert (sizetype, index); - len = fold_convert (sizetype, len); - - tree nc2 = OMP_CLAUSE_CHAIN (nc); - if (nc2 - && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (nc2) - == GOMP_MAP_GRID_STRIDE)) - { - stride = OMP_CLAUSE_DECL (nc2); - stride = fold_convert (sizetype, stride); - if (OMP_CLAUSE_SIZE (nc2)) - { - /* If the element size is not the same as the - distance between two adjacent array - elements (in the innermost dimension), - retrieve the latter value ("span") from the - size field of the stride. We only expect to - see one such field per array. */ - gcc_assert (!span); - span = OMP_CLAUSE_SIZE (nc2); - span = fold_convert (sizetype, span); - } - nc = nc2; - } - - 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 - in libgomp. */ - tree dtype = TYPE_DOMAIN (tdims[tdim]); - tree minval = TYPE_MIN_VALUE (dtype); - tree maxval = TYPE_MAX_VALUE (dtype); - minval = fold_convert (sizetype, minval); - maxval = fold_convert (sizetype, maxval); - dim = size_binop (MINUS_EXPR, maxval, minval); - dim = size_binop (PLUS_EXPR, dim, - size_one_node); - arrsize = size_binop (MULT_EXPR, arrsize, dim); - } - else if (pointer_based && !handled_pointer_section) - { - /* Use the selected array section to determine the - size of the array. */ - tree tmp = size_binop (MULT_EXPR, len, stride); - tmp = size_binop (MINUS_EXPR, tmp, stride); - tmp = size_binop (PLUS_EXPR, tmp, size_one_node); - dim = size_binop (PLUS_EXPR, index, tmp); - arrsize = size_binop (MULT_EXPR, arrsize, dim); - handled_pointer_section = true; - } - else - { - if (pointer_based) - error_at (OMP_CLAUSE_LOCATION (c), - "too many array section specifiers " - "for pointer-based array"); - else - error_at (OMP_CLAUSE_LOCATION (c), - "too many array section specifiers " - "for array"); - dim = index = len = stride = error_mark_node; - } - tdim++; - - c = nc; - } - else - { - /* We have more array dimensions than array section - specifiers. Copy the whole span. */ - tree dtype = TYPE_DOMAIN (tdims[tdim]); - tree minval = TYPE_MIN_VALUE (dtype); - tree maxval = TYPE_MAX_VALUE (dtype); - minval = fold_convert (sizetype, minval); - maxval = fold_convert (sizetype, maxval); - dim = size_binop (MINUS_EXPR, maxval, minval); - dim = size_binop (PLUS_EXPR, dim, size_one_node); - len = dim; - index = minval; - nc = c; - } - - if (TREE_CODE (dim) != INTEGER_CST) - TREE_STATIC (dim_tmp) = 0; - - if (TREE_CODE (index) != INTEGER_CST) - TREE_STATIC (index_tmp) = 0; - - if (TREE_CODE (len) != INTEGER_CST) - TREE_STATIC (len_tmp) = 0; - - if (TREE_CODE (stride) != INTEGER_CST) - TREE_STATIC (stride_tmp) = 0; - - tree cidx = size_int (i); - CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim); - CONSTRUCTOR_APPEND_ELT (vindex, cidx, index); - CONSTRUCTOR_APPEND_ELT (vlen, cidx, len); - CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride); - } - - tree bias = size_zero_node; - tree volume = size_one_node; - tree enclosure = size_one_node; - for (i = dims - 1; i >= 0; i--) - { - tree dim = (*vdim)[i].value; - tree index = (*vindex)[i].value; - tree stride = (*vstride)[i].value; - tree len = (*vlen)[i].value; - - /* For the bias we want, e.g.: - - index[0] * stride[0] * dim[1] * dim[2] - + index[1] * stride[1] * dim[2] - + index[2] * stride[2] - - All multiplied by "span" (or "elsize"). */ - - tree index_stride = size_binop (MULT_EXPR, index, stride); - bias = size_binop (PLUS_EXPR, bias, - size_binop (MULT_EXPR, volume, - index_stride)); - volume = size_binop (MULT_EXPR, volume, dim); - - if (i == 0) - { - tree elems_covered = size_binop (MINUS_EXPR, len, - size_one_node); - elems_covered = size_binop (MULT_EXPR, elems_covered, - stride); - elems_covered = size_binop (PLUS_EXPR, elems_covered, - size_one_node); - enclosure = size_binop (MULT_EXPR, enclosure, - elems_covered); - } - else - enclosure = volume; - } - - /* If we don't have a separate span size, use the element size - instead. */ - if (!span) - span = fold_convert (sizetype, elsize); - - /* The size of a volume enclosing the elements to be - transferred. */ - OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, enclosure, span); - /* And the bias of the first element we will update. */ - OMP_CLAUSE_SIZE (dn) = size_binop (MULT_EXPR, bias, span); - - tree cdim = build_constructor (size_arr_type, vdim); - tree cindex = build_constructor (size_arr_type, vindex); - tree clen = build_constructor (size_arr_type, vlen); - tree cstride = build_constructor (size_arr_type, vstride); - - if (TREE_STATIC (dim_tmp)) - DECL_INITIAL (dim_tmp) = cdim; - else - gimplify_assign (dim_tmp, cdim, &ilist); - - if (TREE_STATIC (index_tmp)) - DECL_INITIAL (index_tmp) = cindex; - else - gimplify_assign (index_tmp, cindex, &ilist); - - if (TREE_STATIC (len_tmp)) - DECL_INITIAL (len_tmp) = clen; - else - gimplify_assign (len_tmp, clen, &ilist); - - if (TREE_STATIC (stride_tmp)) - DECL_INITIAL (stride_tmp) = cstride; - else - gimplify_assign (stride_tmp, cstride, &ilist); - - tree desc_type = TREE_TYPE (desc); - - tree ndims_field = TYPE_FIELDS (desc_type); - tree elemsize_field = DECL_CHAIN (ndims_field); - tree span_field = DECL_CHAIN (elemsize_field); - tree dim_field = DECL_CHAIN (span_field); - tree index_field = DECL_CHAIN (dim_field); - tree len_field = DECL_CHAIN (index_field); - tree stride_field = DECL_CHAIN (len_field); - - vec *v; - vec_alloc (v, 7); - - bool all_static = (TREE_STATIC (dim_tmp) - && TREE_STATIC (index_tmp) - && TREE_STATIC (len_tmp) - && TREE_STATIC (stride_tmp)); - - dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node, - NULL_TREE, NULL_TREE); - dim_tmp = build_fold_addr_expr (dim_tmp); - - /* TODO: we could skip all-zeros index. */ - index_tmp = build4 (ARRAY_REF, sizetype, index_tmp, - size_zero_node, NULL_TREE, NULL_TREE); - index_tmp = build_fold_addr_expr (index_tmp); - - len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node, - NULL_TREE, NULL_TREE); - len_tmp = build_fold_addr_expr (len_tmp); - - /* TODO: we could skip all-ones stride. */ - stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp, - size_zero_node, NULL_TREE, NULL_TREE); - stride_tmp = build_fold_addr_expr (stride_tmp); - - elsize = fold_convert (sizetype, elsize); - tree ndims = size_int (dims); - - CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims); - CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize); - CONSTRUCTOR_APPEND_ELT (v, span_field, span); - CONSTRUCTOR_APPEND_ELT (v, dim_field, dim_tmp); - CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp); - CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp); - CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp); - - tree desc_ctor = build_constructor (desc_type, v); - - if (all_static) - { - TREE_STATIC (desc) = 1; - DECL_INITIAL (desc) = desc_ctor; - } - else - gimplify_assign (desc, desc_ctor, &ilist); - - OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc); - c = oc; + lower_omp_target_grid_desc (c, &ilist); nc = c; } else if (!DECL_P (ovar)) From patchwork Wed Aug 5 20:23:42 2026 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 140679 Return-Path: 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 5B8B14BB24D1 for ; Wed, 5 Aug 2026 20:26:40 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5B8B14BB24D1 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=AtLLex+J X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wr1-x42d.google.com (mail-wr1-x42d.google.com [IPv6:2a00:1450:4864:20::42d]) by sourceware.org (Postfix) with ESMTPS id ABCAF4BAE7EC for ; Wed, 5 Aug 2026 20:25:43 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org ABCAF4BAE7EC 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 ABCAF4BAE7EC Authentication-Results: sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::42d ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961544; cv=none; b=VLQi5dx1g+yG8L9lI8CbkL3B+ky6q35UJVCd8gEU8lSzu4DSAbXJCoTwHUNcBPN0d3zVJaKt6rRHRw5BusAL+pPpcV6wUjGAq+6lG35hUqQauDHEL0FE7JtBprh4tkv5lcXaCWVAH6ZBshH28b5Q5RaZ5yD5XrLjxYfsmQRi2jM= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961544; c=relaxed/simple; bh=/cqwAMkFDkQTTapuxd9bsSnnKxhWGcGNmU3csGBt76A=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=F6SCnOEvA4Cl6M2LBLmc76BvIJINMP2Zv5fo7Ff2l1I7lJ0to/mRt2qGiQgE3naNwAqoAYj/eZlquJkdc+Cfa2OH6uaNUDctV2ZD0c1kf2VACkZrvfemfB/aRjNiCSRmIYTVOXm4NdSMUm2TTYBn+QXaeRmkP0HMHO1dSSzq/mY= 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=AtLLex+J DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org ABCAF4BAE7EC Received: by mail-wr1-x42d.google.com with SMTP id ffacd0b85a97d-4799b3f7c83so861405f8f.2 for ; Wed, 05 Aug 2026 13:25:43 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1785961542; x=1786566342; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to:content-type; bh=6G1pEVGFGa2wdfICI6k/gGOAKqq12913MqTRW0Qt5Gk=; b=AtLLex+JVNbcsfl4jDru3atWadIhFIikwMzA0BTb6TbE1/ZSefom/+2rV/NI4RV/zI cupVWT71Op1wR36XS5AFWPy3vtMjgMaahvfFGMfk16eCYpzAqINWm1DM7rdDj3wL6o2p wazmPz7MZFiY4GfHHDI7NvxneNU7zS7Bl9bsTwTUPu3RdIc/n731Oy7P2fFS1ssd8zoc tQVWR4Zd1b+WcjQ5kor5QYVkVWo+2cbpCZ8xxGzs1koJeUUBG9e+pBt+6ThZjEjZhe5V gkt6aJE1VfdsYhygHCyPxlgoKuY0eOSxfBYyADbP/GCDhVTA2A0LTGyOFzxp+8X+GLc9 /RxQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1785961542; x=1786566342; h=content-transfer-encoding:mime-version:references:in-reply-to :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=6G1pEVGFGa2wdfICI6k/gGOAKqq12913MqTRW0Qt5Gk=; b=YySvnfT3/QDAzqMXM8qHzJMuOg/569IGdsbnrEdqRx2Fssmd77cB8hpzqK+hOuv72s GxEE1EBuGGgHc50Ab65O+e3bC1N6/KPvyTpva863GB6C1Dml0gd2/8MrCwz3VRoVq1ev fq8BQQ3h4yL1eq2r3zQcdeIHmQqXp6AYJZoouikyx8JLLAOo3mywc5V+RDVn0ln/fRgh a+kC7OMlndeO7T0vyIq4mJdINr/qjtgtobHaKEREUPhpyPPhjZ8DVfiGqOOB8cQNLhP3 3I2EenCDGBjD49tgLLPKiWDkvlkNtFXja43p3i8s73DyxlHgynKL0I1H4iDWZnYERz6V PcOQ== X-Gm-Message-State: AOJu0YyNRD5fKBTO+qJoAJySjcY7xGBP/Jq9GvUmcr3BFmk2DotMZiTc obOjyjfK1tc3uwdICmrv0KGQLDwav+3g07HlYzn4fz4G5x1vhrylCFn4rCgE58YIZptHgniA9n9 2+E15 X-Gm-Gg: AR+sD13gIhwPrrdVDEDR/cUpWn9Bqpssn8d9jMQcAQLndIUseL/4mYJhDZSg8e3AApB kCecWEuf9JWQuhYMqRGkEjhYrCNWoFHoUysiNa0K0guZ9HKxuJMOejSfPbcUArHweb7c9onDAX+ ShxwJgHh0OT3V509ZJ8Kun3OYVXseSCwJal/Co+EvwgQo4AaZLrE3qXU1xwXh6APzTb5S+wRPVI 1RINrJ32U4K+WIq0NtNNBGEcnDiItKshzjit0fxCmNisevdhyNxF/j6aVeUKgdjTswqdSnb2E4h aoyPf9WxSdpXO8fbK6j/9gz8eTrp6s1o1mYdrsIlP/ndnTyDcrWKTbHebvpNlvoT8j4FH3rW+lB ogTbz9SMLx1Ba2i5zgVS55+X/aCY7E2RUFttcpBq9V73+UXLdJj5ujJJyD1UzV7HIum9TLcvcSo lTy551Tlmo9G5pJf/oPZqLfPmRtyg3/pmsZt4c+RgxvPWER3FI+VaxYg== X-Received: by 2002:a5d:64c8:0:b0:47f:c62e:9cca with SMTP id ffacd0b85a97d-47fec6333f0mr15892328f8f.22.1785961542337; Wed, 05 Aug 2026 13:25:42 -0700 (PDT) Received: from raptor ([62.108.198.223]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-47ff79b431dsm67029f8f.13.2026.08.05.13.25.41 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 05 Aug 2026 13:25:41 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, Paul-Antoine Arras Subject: [PATCH 4/5] openmp: Generalize lower_omp_target_grid_desc to segments and pointer sections Date: Wed, 5 Aug 2026 22:23:42 +0200 Message-ID: <20260805202343.2868178-5-parras@baylibre.com> X-Mailer: git-send-email 2.53.0 In-Reply-To: <20260805202343.2868178-1-parras@baylibre.com> References: <20260805202343.2868178-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-13.0 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, 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 List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org Extend the grid descriptor lowering to build the noncontiguous-array descriptor for array-shaping casts and array-of-pointers sections, not just plain array sections: - Support multiple segments, each starting where a GRID_DIM clause is marked as selecting through a pointer (OMP_CLAUSE_MAP_GRID_DIM_POINTER), tracking per-segment dimension/pointer counts (__nsegments, __seg_ndims, __seg_nptrs) in the runtime descriptor. - Consume GOMP_MAP_SHAPE_DIM clauses to size dimensions introduced by an array-shaping cast, once the decl's own array type runs out of dimensions. - Stop walking GRID_DIM/GRID_STRIDE/SHAPE_DIM clauses once the last one for the current segment has been consumed, rather than assuming a fixed dimension count. - Accept decayed array parameters (no array type left to consult) by falling back to the last pointer type crossed. - Only mark the runtime descriptor fields TREE_STATIC when every contributing dimension is truly constant, instead of unconditionally. Teach the clause-splicing logic in gimplify.cc about the new GOMP_MAP_SHAPE_DIM clause, and add the "lower" dump scans checking the generated descriptors for the tests added alongside the front-end support. gcc/ChangeLog: * gimplify.cc (omp_group_last): Handle GOMP_MAP_SHAPE_DIM. (gimplify_adjust_omp_clauses): Likewise. * omp-low.cc (omp_noncontig_descriptor_type): Add __nsegments, __seg_ndims and __seg_nptrs fields. (lower_omp_target_grid_desc): Support multiple segments, dimensions introduced by an array-shaping cast, decayed array parameters, and only mark the descriptor fields static when every dimension is constant. (lower_omp_target): Handle GOMP_MAP_SHAPE_DIM alongside GOMP_MAP_GRID_DIM/GOMP_MAP_GRID_STRIDE. gcc/testsuite/ChangeLog: * c-c++-common/gomp/array-section-1.c: Add "lower" dump scans. * c-c++-common/gomp/array-section-2.c: Likewise. * c-c++-common/gomp/array-section-3.c: Likewise. * c-c++-common/gomp/array-section-4.c: Likewise. * c-c++-common/gomp/array-section-5.c: Likewise. * c-c++-common/gomp/array-section-6.c: Likewise. * c-c++-common/gomp/array-section-8.c: Likewise. * c-c++-common/gomp/array-section-9.c: Likewise. --- gcc/gimplify.cc | 18 +- gcc/omp-low.cc | 462 ++++++++++++------ .../c-c++-common/gomp/array-section-1.c | 7 + .../c-c++-common/gomp/array-section-2.c | 7 + .../c-c++-common/gomp/array-section-3.c | 7 + .../c-c++-common/gomp/array-section-4.c | 7 + .../c-c++-common/gomp/array-section-5.c | 7 + .../c-c++-common/gomp/array-section-6.c | 1 + .../c-c++-common/gomp/array-section-8.c | 7 + .../c-c++-common/gomp/array-section-9.c | 7 + 10 files changed, 361 insertions(+), 169 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index b0ecf08cf7f..e4671befc55 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -11608,10 +11608,10 @@ omp_group_last (tree *start_p) case GOMP_MAP_TO_GRID: case GOMP_MAP_FROM_GRID: - while (nc - && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + while (nc && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM - || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE)) + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_SHAPE_DIM)) { grp_last_p = &OMP_CLAUSE_CHAIN (c); c = nc; @@ -16933,7 +16933,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; if (OMP_CLAUSE_SIZE (c) == NULL_TREE && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_DIM - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_STRIDE) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_STRIDE + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_SHAPE_DIM) { /* Sanity check: attach/detach map kinds use the size as a bias, and it's never right to use the decl size for such @@ -17036,11 +17037,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, remove = true; } else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_SHAPE_DIM) { - /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily - an lvalue -- e.g. it might be a constant. So handle it - specially here. */ + /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE/SHAPE_DIM isn't + necessarily an lvalue -- e.g. it might be a constant. So + handle it specially here. */ if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 45f6c701f48..128a0848ad3 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1355,6 +1355,21 @@ omp_noncontig_descriptor_type (location_t loc) TREE_CHAIN (field) = fields; fields = field; + field = build_decl (loc, FIELD_DECL, get_identifier ("__nsegments"), + size_type_node); + TREE_CHAIN (field) = fields; + fields = field; + + field = build_decl (loc, FIELD_DECL, get_identifier ("__seg_ndims"), + ptr_size_type); + TREE_CHAIN (field) = fields; + fields = field; + + field = build_decl (loc, FIELD_DECL, get_identifier ("__seg_nptrs"), + ptr_size_type); + TREE_CHAIN (field) = fields; + fields = field; + finish_builtin_struct (t, "__omp_noncontig_desc_type", fields, ptr_type_node); cached = t; @@ -13723,52 +13738,111 @@ convert_from_firstprivate_int (tree var, tree orig_type, bool is_ref, gimplify_assign (tmp, var, gs); return fold_build1 (VIEW_CONVERT_EXPR, type, tmp); +} /* Build the noncontiguous-array descriptor for the - GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause C (whose GOMP_MAP_TO_PSET + GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause C. The GOMP_MAP_TO_PSET sibling clause holds the descriptor decl and the GOMP_MAP_GRID_DIM/ - GOMP_MAP_GRID_STRIDE clauses that follow describe each dimension). Any - gimplification needed to initialize the descriptor is appended to + GOMP_MAP_GRID_STRIDE/GOMP_MAP_SHAPE_DIM clauses that follow describe each + dimension. Then splice those consumed dimension clauses out of the clause + chain. Any gimplification needed to initialize the descriptor is appended to *ILIST_P. */ static void lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) { tree decl = OMP_CLAUSE_DECL (c); - tree dn = OMP_CLAUSE_CHAIN (c); - gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET); - tree desc = OMP_CLAUSE_DECL (dn); - + tree desc_node = OMP_CLAUSE_CHAIN (c); + gcc_assert (OMP_CLAUSE_CODE (desc_node) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (desc_node) == GOMP_MAP_TO_PSET); + tree desc = OMP_CLAUSE_DECL (desc_node); tree oc, elsize = OMP_CLAUSE_SIZE (c); tree type = TREE_TYPE (decl); - int i, dims = 0; + + /* First, count dimensions and record their types. */ + + int i, dims = 0, segments = 0; tree nc; auto_vec tdims; + auto_vec seg_ndims, seg_nptrs; bool pointer_based = false, handled_pointer_section = false; - tree arrsize = size_one_node; - /* Allow a single (maybe strided) array section if we have a - pointer base. */ if (TREE_CODE (decl) == INDIRECT_REF - && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == POINTER_TYPE)) + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) == POINTER_TYPE)) { + /* Allow a single (maybe strided) array section if we have a + pointer base. */ pointer_based = true; dims = 1; + seg_ndims.safe_push (dims); + segments = 1; } else - /* NOTE: Don't treat (e.g. Fortran, fixed-length) strings as - array types here; array section syntax isn't applicable to - strings. */ - for (tree itype = type; - TREE_CODE (itype) == ARRAY_TYPE - && !TYPE_STRING_FLAG (itype); - itype = TREE_TYPE (itype)) - { - tdims.safe_push (itype); - dims++; - } + { + tree dim_clause = OMP_CLAUSE_CHAIN (desc_node); + gcc_assert (OMP_CLAUSE_CODE (dim_clause) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (dim_clause) == GOMP_MAP_GRID_DIM); + int seg_dims = 0; + + tree ptr_type = NULL_TREE; + tree itype = type; + bool have_itype; + while ((have_itype = (TREE_CODE (itype) == ARRAY_TYPE && + /*array section syntax isn't applicable to strings*/ + !TYPE_STRING_FLAG (itype)) + || TREE_CODE (itype) == POINTER_TYPE) + || dim_clause) + { + /* Once the decl type runs out of dimensions + (e.g. with an array-shaping cast), treat further + dim_clauses as having the type of the last pointer + crossed. */ + tree cur_type = have_itype ? itype : ptr_type; + gcc_assert (cur_type); + + if (dim_clause == NULL_TREE && TREE_CODE (cur_type) == POINTER_TYPE) + { + elsize = TYPE_SIZE_UNIT (cur_type); + break; + } + tdims.safe_push (cur_type); + if (TREE_CODE (cur_type) == POINTER_TYPE) + ptr_type = cur_type; + + if (dim_clause && OMP_CLAUSE_MAP_GRID_DIM_POINTER (dim_clause)) + { + gcc_assert (POINTER_TYPE_P (cur_type)); + seg_ndims.safe_push (seg_dims); + seg_dims = 1; + segments++; + } + else + seg_dims++; + dims++; + + if (dim_clause) + { + dim_clause = OMP_CLAUSE_CHAIN (dim_clause); + while ( + dim_clause + && (OMP_CLAUSE_CODE (dim_clause) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (dim_clause) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (dim_clause) + == GOMP_MAP_SHAPE_DIM))) + dim_clause = OMP_CLAUSE_CHAIN (dim_clause); + + if (dim_clause + && (OMP_CLAUSE_CODE (dim_clause) != OMP_CLAUSE_MAP + || OMP_CLAUSE_MAP_KIND (dim_clause) != GOMP_MAP_GRID_DIM)) + dim_clause = NULL_TREE; + } + + if (have_itype) + itype = TREE_TYPE (itype); + } + seg_ndims.safe_push (seg_dims); + segments++; + } unsigned tdim = 0; @@ -13776,156 +13850,199 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) vec *vindex; vec *vlen; vec *vstride; + vec *vseg_ndims; + vec *vseg_nptrs; vec_alloc (vdim, dims); vec_alloc (vindex, dims); vec_alloc (vlen, dims); vec_alloc (vstride, dims); + vec_alloc (vseg_ndims, segments); + vec_alloc (vseg_nptrs, segments - 1); - tree size_arr_type - = build_array_type_nelts (size_type_node, dims); + tree dim_arr_type = build_array_type_nelts (size_type_node, dims); + tree seg_arr_type = build_array_type_nelts (size_type_node, segments); + tree seg1_arr_type + = build_array_type_nelts (size_type_node, MAX (segments - 1, 1)); - tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim"); + tree dim_tmp = create_tmp_var (dim_arr_type, ".omp_dim"); DECL_NAMELESS (dim_tmp) = 1; TREE_ADDRESSABLE (dim_tmp) = 1; TREE_STATIC (dim_tmp) = 1; - tree index_tmp = create_tmp_var (size_arr_type, ".omp_index"); + tree index_tmp = create_tmp_var (dim_arr_type, ".omp_index"); DECL_NAMELESS (index_tmp) = 1; TREE_ADDRESSABLE (index_tmp) = 1; TREE_STATIC (index_tmp) = 1; - tree len_tmp = create_tmp_var (size_arr_type, ".omp_len"); + tree len_tmp = create_tmp_var (dim_arr_type, ".omp_len"); DECL_NAMELESS (len_tmp) = 1; TREE_ADDRESSABLE (len_tmp) = 1; TREE_STATIC (len_tmp) = 1; - tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride"); + tree stride_tmp = create_tmp_var (dim_arr_type, ".omp_stride"); DECL_NAMELESS (stride_tmp) = 1; TREE_ADDRESSABLE (stride_tmp) = 1; TREE_STATIC (stride_tmp) = 1; + tree seg_ndims_tmp = create_tmp_var (seg_arr_type, ".omp_seg_ndims"); + DECL_NAMELESS (seg_ndims_tmp) = 1; + TREE_ADDRESSABLE (seg_ndims_tmp) = 1; + TREE_STATIC (seg_ndims_tmp) = 1; + tree seg_nptrs_tmp = create_tmp_var (seg1_arr_type, ".omp_seg_nptrs"); + DECL_NAMELESS (seg_nptrs_tmp) = 1; + TREE_ADDRESSABLE (seg_nptrs_tmp) = 1; + TREE_STATIC (seg_nptrs_tmp) = 1; oc = c; - c = dn; + c = desc_node; tree span = NULL_TREE; - for (i = 0; i < dims; i++) + for (int j = 0, idim = 0; j < segments; j++) { - nc = OMP_CLAUSE_CHAIN (c); - tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE, - stride = size_one_node; + tree nptrs = size_int (1); - if (nc - && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM) + for (i = 0; i < seg_ndims[j]; i++, idim++) { - index = OMP_CLAUSE_DECL (nc); - len = OMP_CLAUSE_SIZE (nc); + nc = OMP_CLAUSE_CHAIN (c); + tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE, + stride = size_one_node; - index = fold_convert (sizetype, index); - len = fold_convert (sizetype, len); - - tree nc2 = OMP_CLAUSE_CHAIN (nc); - if (nc2 - && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (nc2) - == GOMP_MAP_GRID_STRIDE)) + if (nc && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM) { - stride = OMP_CLAUSE_DECL (nc2); - stride = fold_convert (sizetype, stride); - if (OMP_CLAUSE_SIZE (nc2)) + index = OMP_CLAUSE_DECL (nc); + len = OMP_CLAUSE_SIZE (nc); + + index = fold_convert (sizetype, index); + len = fold_convert (sizetype, len); + + nptrs = size_binop (MULT_EXPR, nptrs, len); + + tree nc2; + while ((nc2 = OMP_CLAUSE_CHAIN (nc)) + && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_SHAPE_DIM)) { - /* If the element size is not the same as the - distance between two adjacent array - elements (in the innermost dimension), - retrieve the latter value ("span") from the - size field of the stride. We only expect to - see one such field per array. */ - gcc_assert (!span); - span = OMP_CLAUSE_SIZE (nc2); - span = fold_convert (sizetype, span); + if (OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_GRID_STRIDE) + { + stride = OMP_CLAUSE_DECL (nc2); + stride = fold_convert (sizetype, stride); + if (OMP_CLAUSE_SIZE (nc2)) + { + /* If the element size is not the same as the + distance between two adjacent array + elements (in the innermost dimension), + retrieve the latter value ("span") from the + size field of the stride. We only expect + to see one such field per array. + ??? Fortran only, e.g. derived-type array ??? */ + gcc_assert (!span); + span = OMP_CLAUSE_SIZE (nc2); + span = fold_convert (sizetype, span); + } + } + else if (OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_SHAPE_DIM) + dim = OMP_CLAUSE_DECL (nc2); + nc = nc2; } - nc = nc2; - } - if (tdim < tdims.length ()) + if (dim == NULL_TREE) + { + if (tdim < tdims.length () + && TREE_CODE (tdims[tdim]) == ARRAY_TYPE) + { + /* We have an array shape -- use that to find the + total size of the data on the target to look up + in libgomp: + dim = maxval - minval + 1 */ + tree dtype = TYPE_DOMAIN (tdims[tdim]); + tree minval = TYPE_MIN_VALUE (dtype); + tree maxval = TYPE_MAX_VALUE (dtype); + minval = fold_convert (sizetype, minval); + maxval = fold_convert (sizetype, maxval); + dim = size_binop (MINUS_EXPR, maxval, minval); + dim = size_binop (PLUS_EXPR, dim, size_one_node); + } + else if ((pointer_based && !handled_pointer_section) + || TREE_CODE (tdims[tdim]) == POINTER_TYPE) + { + /* Use the selected array section to determine the + size of the array: + dim = index + len * stride - stride + 1 */ + tree tmp = size_binop (MULT_EXPR, len, stride); + tmp = size_binop (MINUS_EXPR, tmp, stride); + tmp = size_binop (PLUS_EXPR, tmp, size_one_node); + dim = size_binop (PLUS_EXPR, index, tmp); + handled_pointer_section = true; + } + else + { + if (pointer_based) + error_at (OMP_CLAUSE_LOCATION (c), + "too many array section specifiers " + "for pointer-based array"); + else + error_at (OMP_CLAUSE_LOCATION (c), + "too many array section specifiers " + "for array"); + dim = index = len = stride = error_mark_node; + } + } + tdim++; + + c = nc; + } + else if (TREE_CODE (tdims[tdim]) == ARRAY_TYPE) { - /* We have an array shape -- use that to find the - total size of the data on the target to look up - in libgomp. */ + /* We have more array dimensions than array section + specifiers. Copy the whole span. */ tree dtype = TYPE_DOMAIN (tdims[tdim]); tree minval = TYPE_MIN_VALUE (dtype); tree maxval = TYPE_MAX_VALUE (dtype); minval = fold_convert (sizetype, minval); maxval = fold_convert (sizetype, maxval); dim = size_binop (MINUS_EXPR, maxval, minval); - dim = size_binop (PLUS_EXPR, dim, - size_one_node); - arrsize = size_binop (MULT_EXPR, arrsize, dim); + dim = size_binop (PLUS_EXPR, dim, size_one_node); + len = dim; + index = minval; + nc = c; } - else if (pointer_based && !handled_pointer_section) - { - /* Use the selected array section to determine the - size of the array. */ - tree tmp = size_binop (MULT_EXPR, len, stride); - tmp = size_binop (MINUS_EXPR, tmp, stride); - tmp = size_binop (PLUS_EXPR, tmp, size_one_node); - dim = size_binop (PLUS_EXPR, index, tmp); - arrsize = size_binop (MULT_EXPR, arrsize, dim); - handled_pointer_section = true; - } - else - { - if (pointer_based) - error_at (OMP_CLAUSE_LOCATION (c), - "too many array section specifiers " - "for pointer-based array"); - else - error_at (OMP_CLAUSE_LOCATION (c), - "too many array section specifiers " - "for array"); - dim = index = len = stride = error_mark_node; - } - tdim++; - c = nc; + if (TREE_CODE (dim) != INTEGER_CST) + TREE_STATIC (dim_tmp) = 0; + + if (TREE_CODE (index) != INTEGER_CST) + TREE_STATIC (index_tmp) = 0; + + if (TREE_CODE (len) != INTEGER_CST) + TREE_STATIC (len_tmp) = 0; + + if (TREE_CODE (stride) != INTEGER_CST) + TREE_STATIC (stride_tmp) = 0; + + tree cidx = size_int (idim); + CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim); + CONSTRUCTOR_APPEND_ELT (vindex, cidx, index); + CONSTRUCTOR_APPEND_ELT (vlen, cidx, len); + CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride); } - else + + tree t = size_int (seg_ndims[j]); + if (TREE_CODE (t) != INTEGER_CST) + TREE_STATIC (seg_ndims_tmp) = 0; + CONSTRUCTOR_APPEND_ELT (vseg_ndims, size_int (j), t); + + if (j < segments - 1) { - /* We have more array dimensions than array section - specifiers. Copy the whole span. */ - tree dtype = TYPE_DOMAIN (tdims[tdim]); - tree minval = TYPE_MIN_VALUE (dtype); - tree maxval = TYPE_MAX_VALUE (dtype); - minval = fold_convert (sizetype, minval); - maxval = fold_convert (sizetype, maxval); - dim = size_binop (MINUS_EXPR, maxval, minval); - dim = size_binop (PLUS_EXPR, dim, size_one_node); - len = dim; - index = minval; - nc = c; + if (TREE_CODE (nptrs) != INTEGER_CST) + TREE_STATIC (seg_nptrs_tmp) = 0; + CONSTRUCTOR_APPEND_ELT (vseg_nptrs, size_int (j), nptrs); } - - if (TREE_CODE (dim) != INTEGER_CST) - TREE_STATIC (dim_tmp) = 0; - - if (TREE_CODE (index) != INTEGER_CST) - TREE_STATIC (index_tmp) = 0; - - if (TREE_CODE (len) != INTEGER_CST) - TREE_STATIC (len_tmp) = 0; - - if (TREE_CODE (stride) != INTEGER_CST) - TREE_STATIC (stride_tmp) = 0; - - tree cidx = size_int (i); - CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim); - CONSTRUCTOR_APPEND_ELT (vindex, cidx, index); - CONSTRUCTOR_APPEND_ELT (vlen, cidx, len); - CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride); } tree bias = size_zero_node; tree volume = size_one_node; tree enclosure = size_one_node; - for (i = dims - 1; i >= 0; i--) + int last_seg_start = dims - seg_ndims[segments - 1]; + for (i = dims - 1; i >= last_seg_start; i--) { tree dim = (*vdim)[i].value; tree index = (*vindex)[i].value; @@ -13942,27 +14059,23 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) tree index_stride = size_binop (MULT_EXPR, index, stride); bias = size_binop (PLUS_EXPR, bias, - size_binop (MULT_EXPR, volume, - index_stride)); + size_binop (MULT_EXPR, volume, index_stride)); volume = size_binop (MULT_EXPR, volume, dim); - if (i == 0) + if (i == last_seg_start) { - tree elems_covered = size_binop (MINUS_EXPR, len, - size_one_node); - elems_covered = size_binop (MULT_EXPR, elems_covered, - stride); - elems_covered = size_binop (PLUS_EXPR, elems_covered, - size_one_node); - enclosure = size_binop (MULT_EXPR, enclosure, - elems_covered); + /* elems_covered = (len - 1) * stride + 1 */ + tree elems_covered = size_binop (MINUS_EXPR, len, size_one_node); + elems_covered = size_binop (MULT_EXPR, elems_covered, stride); + elems_covered = size_binop (PLUS_EXPR, elems_covered, size_one_node); + enclosure = size_binop (MULT_EXPR, enclosure, elems_covered); } else enclosure = volume; } - /* If we don't have a separate span size, use the element size - instead. */ + /* If we don't have a separate span size (??? Fortran only ???), use the + element size instead. */ if (!span) span = fold_convert (sizetype, elsize); @@ -13970,12 +14083,14 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) transferred. */ OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, enclosure, span); /* And the bias of the first element we will update. */ - OMP_CLAUSE_SIZE (dn) = size_binop (MULT_EXPR, bias, span); + OMP_CLAUSE_SIZE (desc_node) = size_binop (MULT_EXPR, bias, span); - tree cdim = build_constructor (size_arr_type, vdim); - tree cindex = build_constructor (size_arr_type, vindex); - tree clen = build_constructor (size_arr_type, vlen); - tree cstride = build_constructor (size_arr_type, vstride); + tree cdim = build_constructor (dim_arr_type, vdim); + tree cindex = build_constructor (dim_arr_type, vindex); + tree clen = build_constructor (dim_arr_type, vlen); + tree cstride = build_constructor (dim_arr_type, vstride); + tree cseg_ndims = build_constructor (seg_arr_type, vseg_ndims); + tree cseg_nptrs = build_constructor (seg1_arr_type, vseg_nptrs); if (TREE_STATIC (dim_tmp)) DECL_INITIAL (dim_tmp) = cdim; @@ -13997,6 +14112,16 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) else gimplify_assign (stride_tmp, cstride, ilist_p); + if (TREE_STATIC (seg_ndims_tmp)) + DECL_INITIAL (seg_ndims_tmp) = cseg_ndims; + else + gimplify_assign (seg_ndims_tmp, cseg_ndims, ilist_p); + + if (TREE_STATIC (seg_nptrs_tmp)) + DECL_INITIAL (seg_nptrs_tmp) = cseg_nptrs; + else + gimplify_assign (seg_nptrs_tmp, cseg_nptrs, ilist_p); + tree desc_type = TREE_TYPE (desc); tree ndims_field = TYPE_FIELDS (desc_type); @@ -14006,35 +14131,47 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) tree index_field = DECL_CHAIN (dim_field); tree len_field = DECL_CHAIN (index_field); tree stride_field = DECL_CHAIN (len_field); + tree nsegments_field = DECL_CHAIN (stride_field); + tree seg_ndims_field = DECL_CHAIN (nsegments_field); + tree seg_nptrs_field = DECL_CHAIN (seg_ndims_field); vec *v; vec_alloc (v, 7); - bool all_static = (TREE_STATIC (dim_tmp) - && TREE_STATIC (index_tmp) - && TREE_STATIC (len_tmp) - && TREE_STATIC (stride_tmp)); + bool all_static + = (TREE_STATIC (dim_tmp) && TREE_STATIC (index_tmp) && TREE_STATIC (len_tmp) + && TREE_STATIC (stride_tmp) && TREE_STATIC (seg_ndims_tmp) + && TREE_STATIC (seg_nptrs_tmp)); - dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node, - NULL_TREE, NULL_TREE); + dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node, NULL_TREE, + NULL_TREE); dim_tmp = build_fold_addr_expr (dim_tmp); /* TODO: we could skip all-zeros index. */ - index_tmp = build4 (ARRAY_REF, sizetype, index_tmp, - size_zero_node, NULL_TREE, NULL_TREE); + index_tmp = build4 (ARRAY_REF, sizetype, index_tmp, size_zero_node, NULL_TREE, + NULL_TREE); index_tmp = build_fold_addr_expr (index_tmp); - len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node, - NULL_TREE, NULL_TREE); + len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node, NULL_TREE, + NULL_TREE); len_tmp = build_fold_addr_expr (len_tmp); /* TODO: we could skip all-ones stride. */ - stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp, - size_zero_node, NULL_TREE, NULL_TREE); + stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp, size_zero_node, + NULL_TREE, NULL_TREE); stride_tmp = build_fold_addr_expr (stride_tmp); + seg_ndims_tmp = build4 (ARRAY_REF, sizetype, seg_ndims_tmp, size_zero_node, + NULL_TREE, NULL_TREE); + seg_ndims_tmp = build_fold_addr_expr (seg_ndims_tmp); + + seg_nptrs_tmp = build4 (ARRAY_REF, sizetype, seg_nptrs_tmp, size_zero_node, + NULL_TREE, NULL_TREE); + seg_nptrs_tmp = build_fold_addr_expr (seg_nptrs_tmp); + elsize = fold_convert (sizetype, elsize); tree ndims = size_int (dims); + tree nsegments = size_int (segments); CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims); CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize); @@ -14043,6 +14180,9 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp); CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp); CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp); + CONSTRUCTOR_APPEND_ELT (v, nsegments_field, nsegments); + CONSTRUCTOR_APPEND_ELT (v, seg_ndims_field, seg_ndims_tmp); + CONSTRUCTOR_APPEND_ELT (v, seg_nptrs_field, seg_nptrs_tmp); tree desc_ctor = build_constructor (desc_type, v); @@ -14054,9 +14194,7 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p) else gimplify_assign (desc, desc_ctor, ilist_p); - OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc); -} - + OMP_CLAUSE_CHAIN (desc_node) = OMP_CLAUSE_CHAIN (nc); } /* Lower the GIMPLE_OMP_TARGET in the current statement @@ -14208,6 +14346,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FROM_GRID: case GOMP_MAP_GRID_DIM: case GOMP_MAP_GRID_STRIDE: + case GOMP_MAP_SHAPE_DIM: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: @@ -14245,7 +14384,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) while ((nc = OMP_CLAUSE_CHAIN (c)) && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM - || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE)) + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_SHAPE_DIM)) c = nc; map_cnt += 2; continue; diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-1.c b/gcc/testsuite/c-c++-common/gomp/array-section-1.c index 99756bbf4c2..e787466d1d1 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-1.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c @@ -13,6 +13,13 @@ void fixed_index (void) int *x[DIM1]; #pragma omp target update to(x[2][ :DIM2]) /* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[2\] = [{]5, 10[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[2\] = [{]2, 0[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[2\] = [{]1, 10[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[2\] = [{]1, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[2\] = [{]1, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_x\.[0-9]+ = [{]\.__ndims=2, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } /* Index, length and stride need not be literal constants -- a variable diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-2.c b/gcc/testsuite/c-c++-common/gomp/array-section-2.c index 15bc2cb28ba..271b5ce31b5 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-2.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-2.c @@ -12,4 +12,11 @@ void two_d_array_of_pointers (void) int *y[DIM1][DIM2]; #pragma omp target update to(y[0][ :2]) /* { dg-final { scan-tree-dump {map\(to_grid:y \[len: [0-9]+\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:0 \[len: 2\]\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[2\] = [{]5, 5[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[2\] = [{]0, 0[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[2\] = [{]1, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[2\] = [{]1, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[1\] = [{]2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[1\] = [{][}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_y\.[0-9]+ = [{]\.__ndims=2, \.__elemsize=8, \.__span=8, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=1, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-3.c b/gcc/testsuite/c-c++-common/gomp/array-section-3.c index f7b17475b5d..f4a35e839db 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-3.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-3.c @@ -16,6 +16,13 @@ void shape_cast_two_fixed_indices_strided_row (void) int *x[DIM1][DIM2]; #pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2][0:2]) /* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[4\] = [{]3, 4, 6, 6[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[4\] = [{]1, 1, 1, 0[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[4\] = [{]1, 1, 3, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[4\] = [{]1, 1, 2, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[2\] = [{]2, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_x\.[0-9]+ = [{]\.__ndims=4, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } void shape_cast_length_one_range_index (void) diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-4.c b/gcc/testsuite/c-c++-common/gomp/array-section-4.c index ce446b931c2..f9a9e8866aa 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-4.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-4.c @@ -18,4 +18,11 @@ void three_consecutive_fixed_indices (void) int *w[D1][D2][D3]; #pragma omp target update to(w[1][2][0][2:4]) /* { dg-final { scan-tree-dump {map\(to_grid:w \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 4\] \[pointer\]\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[4\] = [{]3, 3, 4, 6[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[4\] = [{]1, 2, 0, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[4\] = [{]1, 1, 1, 4[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[4\] = [{]1, 1, 1, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[2\] = [{]3, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_w\.[0-9]+ = [{]\.__ndims=4, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-5.c b/gcc/testsuite/c-c++-common/gomp/array-section-5.c index 9bd8a251393..bd136ead018 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-5.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-5.c @@ -15,6 +15,13 @@ void shape_cast_omitted_column_dim_strided (void) int *x[DIM1][DIM2]; #pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2]) /* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 6\]\) map\(shape_dim:6\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[4\] = [{]3, 4, 6, 6[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[4\] = [{]1, 1, 1, 0[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[4\] = [{]1, 1, 3, 6[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[4\] = [{]1, 1, 2, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[2\] = [{]2, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_x\.[0-9]+ = [{]\.__ndims=4, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } void shape_cast_omitted_column_dim_length_one_range (void) diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-6.c b/gcc/testsuite/c-c++-common/gomp/array-section-6.c index e90f4e66188..0199f725eda 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-6.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-6.c @@ -15,4 +15,5 @@ void plain_pointer_shape_cast (void) #pragma omp target update to((([N]) ptr)[10:30]) /* { dg-final { scan-tree-dump {to\(VIEW_CONVERT_EXPR\(\*ptr\)\[10\] \[len: 120\]\)} "original" } } */ /* { dg-final { scan-tree-dump-not "to_grid" "original" } } */ +/* { dg-final { scan-tree-dump-not "__ndims" "lower" } } */ } diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-8.c b/gcc/testsuite/c-c++-common/gomp/array-section-8.c index 4873138362a..986d774782f 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-8.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-8.c @@ -12,6 +12,13 @@ void three_segment_one_dim_per_segment (void) int **arr[DIM1]; #pragma omp target update from(arr[2][0:2:2][3:5:2]) /* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 5\] \[pointer\]\) map\(grid_stride:2\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[3\] = [{]4, 3, 12[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[3\] = [{]2, 0, 3[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[3\] = [{]1, 2, 5[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[3\] = [{]1, 2, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[3\] = [{]1, 1, 1[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[2\] = [{]1, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_arr\.[0-9]+ = [{]\.__ndims=3, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=3, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } /* Same shape, but every index/length/stride across all three segments is diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-9.c b/gcc/testsuite/c-c++-common/gomp/array-section-9.c index 9d2e0ceae09..eeecb1faa57 100644 --- a/gcc/testsuite/c-c++-common/gomp/array-section-9.c +++ b/gcc/testsuite/c-c++-common/gomp/array-section-9.c @@ -18,4 +18,11 @@ void three_segment_two_dims_per_segment (void) p1_t arr[DIM1A][DIM1B]; #pragma omp target update from(arr[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2]) /* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:0 \[len: 2\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_dim:3 \[len: 5\]\) map\(grid_stride:2\)} "original" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_dim\.[0-9]+\[6\] = [{]2, 3, 3, 3, 2, 20[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_index\.[0-9]+\[6\] = [{]1, 0, 0, 0, 0, 3[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_len\.[0-9]+\[6\] = [{]1, 2, 2, 2, 2, 5[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_stride\.[0-9]+\[6\] = [{]1, 2, 2, 2, 1, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_ndims\.[0-9]+\[3\] = [{]2, 2, 2[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static long unsigned int \.omp_seg_nptrs\.[0-9]+\[2\] = [{]2, 4[}];} "lower" } } */ +/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_arr\.[0-9]+ = [{]\.__ndims=6, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\], \.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\], \.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=3, \.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\], \.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */ } From patchwork Wed Aug 5 20:23:43 2026 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul-Antoine Arras X-Patchwork-Id: 140683 Return-Path: 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 DDDEE4BB24CB for ; Wed, 5 Aug 2026 20:29:12 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org DDDEE4BB24CB 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=phd88CGk X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x332.google.com (mail-wm1-x332.google.com [IPv6:2a00:1450:4864:20::332]) by sourceware.org (Postfix) with ESMTPS id 241A34BB1C3D for ; Wed, 5 Aug 2026 20:25:45 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 241A34BB1C3D 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 241A34BB1C3D Authentication-Results: sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::332 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961553; cv=none; b=Cir8pzf/WmDJI/M2z/dxIaldr9z8Mh+ie41vBgrGUt4xXZlQPUfeF/Jq1kUguA0fuzRqVPUoHKc0aQs3YrVcNUNJ8JBoPUroRnwV8JByESOkDzqlIMKMbyUvIqA7LoJ5AqnpVdRKyYY5lZ6x+ZCRq+PkDMzVbekMLs30mU0ba98= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1785961553; c=relaxed/simple; bh=wXICT0qADc38/iVg18Sykva7CQLzNrQC3/315hA3oM8=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=RZOYvgxJWiz1i3dEZT21zeeoY2RcHVpJ8soBoJp60F/42bD5bePdLUs46uUGQG8UVV8RisLAUG5LX+CcGrO13NsGeZZgKvBe/jzEcTKi1V0azyxlbcw1szdV8oIfZlBQ2hlG7Yv19lB89yuLu7rVoMxREZPsXSvezaqVcnWGQKg= 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=phd88CGk DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 241A34BB1C3D Received: by mail-wm1-x332.google.com with SMTP id 5b1f17b1804b1-4980fe6b3beso1936065e9.0 for ; Wed, 05 Aug 2026 13:25:45 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1785961544; x=1786566344; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to:content-type; bh=3bGlqBLd3EGsnS63VBvhAJNEOj1IDm67GwHwVRBByxM=; b=phd88CGkAnS2PDTKUKZHElUHvRf+hT2JMeQ4ZoS5PCmfruuIcfLBR5bNRqCeU0NVwp j1hcJQiYj0iimjP+hH9RdlIc4hP7jjeMNSO7mcDglnX2pnifQVW3HHtk6cRJbjhTMq7U 3SyKxIMtv0Y3o4Ri8hRaTcEtxeLnhYYrMUZ9Tx+brHBIAbir9gFZu6tiMCxmGmpvQwuh dacrrOo3wgT1uJyJfC3cDTw7yt3nPn7iZ1HQz0aO+ed3n97iueOWtqRVqSn4l8iA+ySO iix6cg5Fa9pKtR4qkx/XPNvrRp6jbDplh12MaWB3iVKtOMI7O5LfAxiVP1JvW7RbtzVa FUQw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1785961544; x=1786566344; h=content-transfer-encoding:mime-version:references:in-reply-to :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=3bGlqBLd3EGsnS63VBvhAJNEOj1IDm67GwHwVRBByxM=; b=LMbl4Nb/jxWfQqbzFeiWKrhDQyn075cw/rylLbIq2ft7T4h2nsUupy6tZqBltT6P93 rg/moCokJD7117rGQVb3412OTh6CAaWvcg65OrEkCFBjsqPLGKM7rvlTxOOUoQ5TD6eZ xM5kFwZkf5uiOvyPa8u1mOdKn+hmWO9OMJvubBZW/UjDAYc+MhJrUv4XsPrPxSOrC/Xu rApB1UWkQYl7QcIb1Ihx96a5VeUqgJsMbhSPbxONMqcRJ1YrvFff2O2loUwO/Rxr2bTU eGO57tWIRPXQAbZ3yXGoV8zVwWx7n3t+jKGUnELRXEfd9pnFuT6CS2t1kMamaH8skCiQ tO0A== X-Gm-Message-State: AOJu0YxpGzZ6YK5XHd1S5g03JSCApkQWTIIsqMGHFkdPz36FLZplXws8 3A7SmBTPuzhPxMWD7x5yj8n9ozb84ZPNbdiOHHqcqck//xlaLqrwaKMYLvMIeWfqJRMpAUm5V1+ W0YY5 X-Gm-Gg: AR+sD11IsXo+5MXnbAl7AMvc45HnJYDw9yYokaCS3MdTHr8gwbhDzKhBL7zp+x55kIT gcAfs2eMANmBzcYlBQ6EFi1Q8gpMX4+wlDdvZ5q1QPEuV4r7h7A7uN6bP9hFqN+jq9GDCllFVPz PwJIpKpdnrvKnbiQYs4ZH2BSW30hCkafJ85qmLkZeFtBvVUTTgkYvcPX51uwJD30aiUXnSWaeCQ /vLuc7c86GIo0LBuVMW1ujj4I2YvVmlqMqycusWDibn3xnh8ELhH6qboIyVAwaXmhkA4QrMctZS Z6fV54HLcfZqA/iyIPUS0UuxooLHEvnfTWOPIxnCvh+hut9foYvPi+I/SNx1B8paKGsz++74cTF /DKYQoSEYEE5LBp80v52k1DSM/jIvutc9Gyqq1v3JqS2FzqvcwEZQvUUr6rEaTUcXcWFmXkBEWr L4dfJ7UwyEkY2hdv24lp2zHk7ylSW+ezD+s21MHQqg225hrbFDbowkQQ== X-Received: by 2002:a05:600c:468c:b0:499:4d4b:3719 with SMTP id 5b1f17b1804b1-49952ef3dfemr19427225e9.4.1785961543491; Wed, 05 Aug 2026 13:25:43 -0700 (PDT) Received: from raptor ([62.108.198.223]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-47ff79b431dsm67029f8f.13.2026.08.05.13.25.42 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 05 Aug 2026 13:25:42 -0700 (PDT) From: Paul-Antoine Arras To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, Paul-Antoine Arras Subject: [PATCH 5/5] openmp: Support multi-segment noncontiguous descriptors in libgomp Date: Wed, 5 Aug 2026 22:23:43 +0200 Message-ID: <20260805202343.2868178-6-parras@baylibre.com> X-Mailer: git-send-email 2.53.0 In-Reply-To: <20260805202343.2868178-1-parras@baylibre.com> References: <20260805202343.2868178-1-parras@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-13.1 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, 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 List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~patchwork=sourceware.org@gcc.gnu.org Extend the runtime noncontiguous-array descriptor and its target-side gather/scatter loops in target.c to walk multiple segments (each crossing a further pointer indirection) instead of assuming a single flat set of dimensions, driven by the new nsegments/seg_ndims/seg_nptrs fields. Add runtime tests exercising array-shaping casts and array-of-pointers sections with one, two, and three segments in C and C++. libgomp/ChangeLog: * libgomp.h (omp_noncontig_array_desc): Add nsegments, seg_ndims and seg_nptrs fields; document all fields. * target.c (omp_target_update_segment): New, extracted and generalized to walk a single segment of a noncontiguous-array descriptor. (gomp_update): Call it once per segment instead of assuming a single flat set of dimensions. * testsuite/libgomp.c++/array-section-1.C: New test. * testsuite/libgomp.c++/array-section-2.C: New test. * testsuite/libgomp.c-c++-common/array-section-1.c: New test. * testsuite/libgomp.c-c++-common/array-section-10.c: New test. * testsuite/libgomp.c-c++-common/array-section-11.c: New test. * testsuite/libgomp.c-c++-common/array-section-2.c: New test. * testsuite/libgomp.c-c++-common/array-section-3.c: New test. * testsuite/libgomp.c-c++-common/array-section-4.c: New test. * testsuite/libgomp.c-c++-common/array-section-5.c: New test. * testsuite/libgomp.c-c++-common/array-section-6.c: New test. * testsuite/libgomp.c-c++-common/array-section-7.c: New test. * testsuite/libgomp.c-c++-common/array-section-8.c: New test. * testsuite/libgomp.c-c++-common/array-section-9.c: New test. --- 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 +++++++++++ 15 files changed, 1686 insertions(+), 37 deletions(-) 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 diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 9f13f9e32c2..eadecbc0055 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1385,13 +1385,23 @@ struct target_mem_desc { omp-low.cc:omp_noncontig_descriptor_type. */ typedef struct { - size_t ndims; - size_t elemsize; - size_t span; - size_t *dim; - size_t *index; - size_t *length; - size_t *stride; + size_t ndims; /* Number of array dimensions */ + size_t elemsize; /* Byte size of each innermost element */ + size_t + span; /* Byte distance between two logically-adjacent elements. Defaults to + elemsize but can be more if stride is specified. */ + size_t *dim; /* Per-dimension total extent */ + size_t *index; /* Per-dimension start offset */ + size_t *length; /* Per-dimension count of selected elements (1 for a fixed + index, more for a range) */ + size_t *stride; /* Per-dimension element stride (defaults to 1) */ + size_t nsegments; /* Number of sets of contiguous dimensions + (= 1 + number of indirections) */ + size_t *seg_ndims; /* Per-segment dimension count */ + size_t *seg_nptrs; /* For each segment junction, number of pointers + selected by the indirection between segment k and k+1 + (1 for a fixed index, N for a strided/ranged dimension + in segment k) */ } omp_noncontig_array_desc; diff --git a/libgomp/target.c b/libgomp/target.c index 859e2de1847..6c5650d7f42 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2693,6 +2693,103 @@ omp_target_memcpy_rect_worker (void *, const void *, size_t, size_t, int, struct gomp_device_descr *, size_t *tmp_size, void **tmp); +/* Transfer BASE to or from the device for one segment of DESC's pointer- + indirected dimensions, recursing through each indirection until the + innermost (contiguous) segment, where the actual rectangular memcpy + happens. SEG_START holds each segment's starting index into DESC's flat + per-dimension arrays (dim/index/length/stride). SEG is the segment + currently being processed. BASE is the host address this segment + indexes into -- the original host pointer for segment 0, or the pointee + reached through the previous segment's selected pointer otherwise. */ + +static void +omp_target_update_segment (omp_noncontig_array_desc *desc, + struct gomp_device_descr *devicep, bool to_not_from, + size_t seg_start[], int seg, void *base) +{ + size_t start = seg_start[seg]; + size_t ndims = desc->seg_ndims[seg]; + + /* Inner segment: do the actual memory transfer. */ + if (seg == desc->nsegments - 1) + { + /* Get device address. */ + struct splay_tree_key_s cur_node; + cur_node.host_start = (uintptr_t) base; + cur_node.host_end = cur_node.host_start + 1; + splay_tree_key tn = splay_tree_lookup (&devicep->mem_map, &cur_node); + if (tn == NULL) + { + gomp_mutex_unlock (&devicep->lock); + gomp_error ("pointer not mapped"); + return; + } + + /* Do rectangular memcpy. */ + void *devaddr = (void *) (tn->tgt->tgt_start + tn->tgt_offset); + size_t inner_seg = desc->nsegments - 1; + size_t inner_dim = desc->ndims - desc->seg_ndims[inner_seg]; + size_t tmp_size = 0; + void *tmp = NULL; + if (to_not_from) + omp_target_memcpy_rect_worker ( + devaddr, base, desc->elemsize, desc->span, + desc->seg_ndims[desc->nsegments - 1], &desc->length[inner_dim], + &desc->stride[inner_dim], &desc->index[inner_dim], + &desc->index[inner_dim], &desc->dim[inner_dim], &desc->dim[inner_dim], + devicep, NULL, &tmp_size, &tmp); + else + omp_target_memcpy_rect_worker ( + base, devaddr, desc->elemsize, desc->span, + desc->seg_ndims[desc->nsegments - 1], &desc->length[inner_dim], + &desc->stride[inner_dim], &desc->index[inner_dim], + &desc->index[inner_dim], &desc->dim[inner_dim], &desc->dim[inner_dim], + NULL, devicep, &tmp_size, &tmp); + return; + } + + /* Intermediate segment: compute position of each pointer and thread pointee + through recursion. */ + for (size_t j = 0; j < desc->seg_nptrs[seg]; j++) + { + size_t rem = j; + size_t idx[ndims]; + for (size_t dd = 0; dd < ndims; dd++) + { + size_t d = start + ndims - 1 - dd; + size_t kk = rem % desc->length[d]; + rem /= desc->length[d]; + idx[d - start] = desc->index[d] + kk * desc->stride[d]; + } + size_t pos = 0; + for (size_t d = start; d < start + ndims; d++) + pos = pos * desc->dim[d] + idx[d - start]; + + void **host_ptr = (void **) base + pos; + void *inner_base = *host_ptr; + + omp_target_update_segment (desc, devicep, to_not_from, seg_start, seg + 1, + inner_base); + } +} + +/* Perform a "target update" on DEVICEP, copying each of the MAPNUM regions + described by HOSTADDRS/SIZES/KINDS (KINDS encoded per SHORT_MAPKIND) + between host and device, using whichever device mapping is already + recorded in DEVICEP's memory map. A region not currently mapped is + silently skipped, unless its map kind requires the region to be present, + in which case that is a fatal error. + + A GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID entry describes a non-contiguous + (reshaped or strided) array section rather than a single flat region; + HOSTADDRS[I + 1] then points to the omp_noncontig_array_desc describing + its dimensions, which is copied one dimension at a time via + omp_target_memcpy_rect_worker and consumes the following entry as well. + + For an ordinary region whose underlying object currently has pointers + being attached, only the individual pointer-sized slots not in the + middle of being attached are updated. */ + static void gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, size_t *sizes, void *kinds, bool short_mapkind) @@ -2729,39 +2826,58 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, omp_noncontig_array_desc *desc = (omp_noncontig_array_desc *) hostaddrs[i + 1]; size_t bias = sizes[i + 1]; - cur_node.host_start = (uintptr_t) hostaddrs[i] + bias; - cur_node.host_end = cur_node.host_start + sizes[i]; - splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node); - if (n) + + if (desc->nsegments == 1) { - if (n->aux && n->aux->attach_count) + cur_node.host_start = (uintptr_t) hostaddrs[i] + bias; + cur_node.host_end = cur_node.host_start + sizes[i]; + splay_tree_key n + = splay_tree_lookup (&devicep->mem_map, &cur_node); + if (n) { - gomp_mutex_unlock (&devicep->lock); - gomp_error ("noncontiguous update with attached pointers"); - return; + if (n->aux && n->aux->attach_count) + { + gomp_mutex_unlock (&devicep->lock); + gomp_error ( + "noncontiguous update with attached pointers"); + return; + } + void *devaddr + = (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start - n->host_start - bias); + size_t tmp_size = 0; + void *tmp = NULL; + if ((kind & typemask) == GOMP_MAP_TO_GRID) + omp_target_memcpy_rect_worker (devaddr, hostaddrs[i], + desc->elemsize, desc->span, + desc->ndims, desc->length, + desc->stride, desc->index, + desc->index, desc->dim, + desc->dim, devicep, NULL, + &tmp_size, &tmp); + else + omp_target_memcpy_rect_worker (hostaddrs[i], devaddr, + desc->elemsize, desc->span, + desc->ndims, desc->length, + desc->stride, desc->index, + desc->index, desc->dim, + desc->dim, NULL, devicep, + &tmp_size, &tmp); } - void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset - + cur_node.host_start - - n->host_start - - bias); - size_t tmp_size = 0; - void *tmp = NULL; - if ((kind & typemask) == GOMP_MAP_TO_GRID) - omp_target_memcpy_rect_worker (devaddr, hostaddrs[i], - desc->elemsize, desc->span, - desc->ndims, desc->length, - desc->stride, desc->index, - desc->index, desc->dim, - desc->dim, devicep, - NULL, &tmp_size, &tmp); - else - omp_target_memcpy_rect_worker (hostaddrs[i], devaddr, - desc->elemsize, desc->span, - desc->ndims, desc->length, - desc->stride, desc->index, - desc->index, desc->dim, - desc->dim, NULL, - devicep, &tmp_size, &tmp); + } + else + { + assert (desc->nsegments != 0); + + /* Compute the start dimension for each segment. */ + size_t seg_start[desc->nsegments]; + seg_start[0] = 0; + for (int s = 1; s < desc->nsegments; s++) + seg_start[s] = seg_start[s - 1] + desc->seg_ndims[s - 1]; + + omp_target_update_segment (desc, devicep, + (kind & typemask) == GOMP_MAP_TO_GRID, + seg_start, 0, hostaddrs[i]); } i++; } diff --git a/libgomp/testsuite/libgomp.c++/array-section-1.C b/libgomp/testsuite/libgomp.c++/array-section-1.C new file mode 100644 index 00000000000..cdd396efbd8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-section-1.C @@ -0,0 +1,204 @@ +// { dg-do run { target offload_device_nonshared_as } } + +/* Noncontiguous "target update" through an array of pointers accessed via + a C++ reference + array-shaped variant. */ + +#include +#include +#include + +#define DIM1 5 +#define DIM2 10 +#define ROWS 6 +#define COLS 6 + +void +test_fixed_index () +{ + int *x[DIM1]; + int *(&x_ref)[DIM1] = x; + int i, j; + + for (i = 0; i < DIM1; i++) + x[i] = (int *) calloc (DIM2, sizeof (int)); + +#pragma omp target enter data map(alloc: x_ref[0:DIM1]) +#pragma omp target enter data map(to: x_ref[2][:DIM2]) + + for (j = 0; j < DIM2; j++) + x_ref[2][j] = j + 1; + +#pragma omp target update to(x_ref[2][:DIM2]) + +#pragma omp target map(alloc: x_ref) + { + for (int j = 0; j < DIM2; j++) + x_ref[2][j] += 1000; + } + + memset (x[2], 0, DIM2 * sizeof (int)); + +#pragma omp target exit data map(from: x_ref[2][:DIM2]) +#pragma omp target exit data map(release: x_ref[0:DIM1]) + + for (j = 0; j < DIM2; j++) + assert (x[2][j] == j + 1 + 1000); + + for (i = 0; i < DIM1; i++) + free (x[i]); +} + +void +test_strided_range () +{ + int *x[DIM1]; + int *(&x_ref)[DIM1] = x; + int i, j; + + for (i = 0; i < DIM1; i++) + x[i] = (int *) calloc (DIM2, sizeof (int)); + +#pragma omp target enter data map(alloc: x_ref[0:DIM1]) + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x_ref[i][:DIM2]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i % 2 == 0) + x_ref[i][j] = 10 * i + j; + +#pragma omp target update to(x_ref[0:3:2][:DIM2]) + +#pragma omp target map(alloc: x_ref) + { + for (int i = 0; i < DIM1; i += 2) + for (int j = 0; j < DIM2; j++) + x_ref[i][j] += 1000; + } + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(from: x_ref[i][:DIM2]) + } + +#pragma omp target exit data map(release: x_ref[0:DIM1]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i % 2 == 0) + assert (x[i][j] == 10 * i + j + 1000); + else + assert (x[i][j] == 0); + + for (i = 0; i < DIM1; i++) + free (x[i]); +} + +void +test_update_from () +{ + int *x[DIM1]; + int *(&x_ref)[DIM1] = x; + int i, j; + + for (i = 0; i < DIM1; i++) + x[i] = (int *) calloc (DIM2, sizeof (int)); + +#pragma omp target enter data map(alloc: x_ref[0:DIM1]) +#pragma omp target enter data map(to: x_ref[3][:DIM2]) + + for (j = 0; j < DIM2; j++) + x_ref[3][j] = 100 + j; + +#pragma omp target update to(x_ref[3][:DIM2]) + +#pragma omp target map(alloc: x_ref) + { + for (int j = 0; j < DIM2; j++) + x_ref[3][j] += 1000; + } + + for (j = 0; j < DIM2; j++) + x_ref[3][j] = -1; + +#pragma omp target update from(x_ref[3][:DIM2]) + + for (j = 0; j < DIM2; j++) + assert (x[3][j] == 100 + j + 1000); + +#pragma omp target exit data map(delete: x_ref[3][:DIM2]) +#pragma omp target exit data map(release: x_ref[0:DIM1]) + + for (i = 0; i < DIM1; i++) + free (x[i]); +} + +void +test_shape_cast () +{ + int *x[DIM1]; + int *(&x_ref)[DIM1] = x; + int i, r, c; + + for (i = 0; i < DIM1; i++) + x[i] = (int *) calloc (ROWS * COLS, sizeof (int)); + +#pragma omp target enter data map(alloc: x_ref[0:DIM1]) + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x_ref[i][:ROWS*COLS]) + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x_ref[1][r * COLS + c] = 100 * r + c; + +#pragma omp target update to((([ROWS][COLS]) x_ref[1])[1:3:2][0:2]) + +#pragma omp target map(alloc: x_ref) + { + for (int r = 1; r <= 5; r += 2) + for (int c = 0; c < 2; c++) + x_ref[1][r * COLS + c] += 1000; + } + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, ROWS * COLS * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(from: x_ref[i][:ROWS*COLS]) + } + +#pragma omp target exit data map(release: x_ref[0:DIM1]) + + for (i = 0; i < DIM1; i++) + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + { + int expected = 0; + if (i == 1 && (r == 1 || r == 3 || r == 5) && c < 2) + expected = 100 * r + c + 1000; + assert (x[i][r * COLS + c] == expected); + } + + for (i = 0; i < DIM1; i++) + free (x[i]); +} + +int +main () +{ + test_fixed_index (); + test_strided_range (); + test_update_from (); + test_shape_cast (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-section-2.C b/libgomp/testsuite/libgomp.c++/array-section-2.C new file mode 100644 index 00000000000..c37f3200cce --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-section-2.C @@ -0,0 +1,130 @@ +// { dg-do run { target offload_device_nonshared_as } } + +/* Test noncontiguous "target update" with three segments (two pointer indirections) + and two dimensions per segment, along with C++ templates and references. */ + +#include +#include + +#define DIM1A 2 +#define DIM1B 3 +#define ROWS1 3 +#define NCOLS 3 +#define ROWS2 3 +#define ROWLEN 20 + +template +static T +encode (int i, int j, int r1, int c1, int r2, int k) +{ + return (T) (i * 1000000 + j * 100000 + r1 * 10000 + c1 * 1000 + + r2 * 100 + k); +} + +template +void +test () +{ + typedef T (*p2_t)[ROWLEN]; + typedef p2_t (*p1_t)[NCOLS]; + + p1_t arr[DIM1A][DIM1B]; + p1_t (&arr_ref)[DIM1A][DIM1B] = arr; + int i, j, r1, c1, r2, k; + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + { + arr_ref[i][j] = (p1_t) malloc (ROWS1 * NCOLS * sizeof (p2_t)); + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + arr_ref[i][j][r1][c1] + = (p2_t) malloc (ROWS2 * ROWLEN * sizeof (T)); + } + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + for (r2 = 0; r2 < ROWS2; r2++) + for (k = 0; k < ROWLEN; k++) + arr_ref[i][j][r1][c1][r2][k] = encode (i, j, r1, c1, r2, k); + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + { +#pragma omp target enter data map(to: arr_ref[i][j][r1][c1][:ROWS2]) + } + + /* Corrupt the whole of ARR_REF[1]'s data, so that precisely matching + the six-dimensional touched set below after the update proves each + dimension of each segment was walked correctly. */ + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + for (r2 = 0; r2 < ROWS2; r2++) + for (k = 0; k < ROWLEN; k++) + arr_ref[1][j][r1][c1][r2][k] = (T) -1; + +#pragma omp target update from(arr_ref[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2]) + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + for (r2 = 0; r2 < ROWS2; r2++) + for (k = 0; k < ROWLEN; k++) + { + T expected; + if (i == 1) + { + int j_touched = (j == 0 || j == 2); + int r1_touched = (r1 == 0 || r1 == 2); + int c1_touched = (c1 == 0 || c1 == 2); + int r2_touched = (r2 == 0 || r2 == 1); + int k_touched + = (k == 3 || k == 5 || k == 7 || k == 9 || k == 11); + if (j_touched && r1_touched && c1_touched + && r2_touched && k_touched) + /* Restored from the device -- proves this exact + combination of dimensions across all three + segments was selected. */ + expected = encode (i, j, r1, c1, r2, k); + else + /* Still corrupted -- proves the update did not + touch anything outside the requested set. */ + expected = (T) -1; + } + else + /* Untouched by either the corruption or the update. */ + expected = encode (i, j, r1, c1, r2, k); + assert (arr_ref[i][j][r1][c1][r2][k] == expected); + } + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + { +#pragma omp target exit data map(delete: arr_ref[i][j][r1][c1][:ROWS2]) + } + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + { + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + free (arr_ref[i][j][r1][c1]); + free (arr_ref[i][j]); + } +} + +int +main () +{ + test (); + test (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-1.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-1.c new file mode 100644 index 00000000000..514367f16a5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-1.c @@ -0,0 +1,181 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Noncontiguous "target update" through one level of pointer indirection, + i.e. an array of dynamically-allocated pointers. */ + +#include +#include +#include + +#define DIM1 5 +#define DIM2 10 + +int main() +{ + int *x[DIM1]; + int i, j; + + for (i = 0; i < DIM1; i++) + x[i] = (int *)calloc (DIM2, sizeof (int)); + +#pragma omp target enter data map(alloc: x[0:DIM1]) + + /* Case 1: Fixed-index update to device through one pointer. */ + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:DIM2]) + } + + for (j = 0; j < DIM2; j++) + x[2][j] = j + 1; + +#pragma omp target update to(x[2][:DIM2]) + +#pragma omp target map(alloc: x) + { + for (int j = 0; j < DIM2; j++) + x[2][j] += 1000; + } + + /* Erase the host copies before pulling back from the device, so the + assertions below only see what actually made it across. */ + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(from: x[i][:DIM2]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i == 2) + assert (x[i][j] == j + 1 + 1000); + else + assert (x[i][j] == 0); + + /* Case 2: Range of pointers: a single directive walks through several + selected pointers, each contributing its own contiguous run. */ + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:DIM2]) + } + + for (i = 1; i <= 3; i++) + for (j = 0; j < DIM2; j++) + x[i][j] = 10 * i + j; + +#pragma omp target update to(x[1:3][:DIM2]) + +#pragma omp target map(alloc: x) + { + for (int i = 1; i <= 3; i++) + for (int j = 0; j < DIM2; j++) + x[i][j] += 1000; + } + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(from: x[i][:DIM2]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i >= 1 && i <= 3) + assert (x[i][j] == 10 * i + j + 1000); + else + assert (x[i][j] == 0); + + /* Case 3: Strided range of pointers. */ + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:DIM2]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i % 2 == 0) + x[i][j] = 10 * i + j; + +#pragma omp target update to(x[0:3:2][:DIM2]) + +#pragma omp target map(alloc: x) + { + for (int i = 0; i < DIM1; i += 2) + for (int j = 0; j < DIM2; j++) + x[i][j] += 1000; + } + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(from: x[i][:DIM2]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i%2 == 0) + assert (x[i][j] == 10 * i + j + 1000); + else + assert (x[i][j] == 0); + + /* Case 4: Update from device through the indirection + target region + computation. */ + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, DIM2 * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:DIM2]) + } + + for (j = 0; j < DIM2; j++) + x[3][j] = 100 + j; + +#pragma omp target update to(x[3][:DIM2]) + +#pragma omp target map(alloc: x) + { + for (int j = 0; j < DIM2; j++) + x[3][j] += 1000; + } + + for (j = 0; j < DIM2; j++) + x[3][j] = -1; + +#pragma omp target update from(x[3][:DIM2]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i == 3) + assert (x[i][j] == 100 + j + 1000); + else + assert (x[i][j] == 0); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(delete: x[i][:DIM2]) + } + +#pragma omp target exit data map(release: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + free (x[i]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-10.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-10.c new file mode 100644 index 00000000000..531f3a641ca --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-10.c @@ -0,0 +1,69 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Check that data transferred through noncontiguous target update can be + accessed from the device. */ + +#include +#include + +#define DIM1 5 +#define DIM2 10 + +int +main () +{ + int *x[DIM1]; + int i, j; + + for (i = 0; i < DIM1; i++) + x[i] = (int *) calloc (DIM2, sizeof (int)); + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + x[i][j] = 100 * i + j; + +#pragma omp target enter data map(alloc: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:DIM2]) + } + + for (j = 0; j < DIM2; j += 2) + x[1][j] = 1000 + j; + +#pragma omp target update to(x[1][0:DIM2/2:2]) + + /* Corrupt the host copy of the whole row, so the assertions at the + end can only pass if the values genuinely came from the device -- + both from the strided update above and from the target region + below. */ + for (j = 0; j < DIM2; j++) + x[1][j] = -1; + +#pragma omp target map(alloc: x) + { + for (int j = 0; j < DIM2; j++) + x[1][j] += 1; + } + +#pragma omp target update from(x[1][:DIM2]) + + for (j = 0; j < DIM2; j++) + { + int expected = (j % 2 == 0) ? (1000 + j + 1) : (100 * 1 + j + 1); + assert (x[1][j] == expected); + } + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(delete: x[i][:DIM2]) + } + +#pragma omp target exit data map(delete: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + free (x[i]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-11.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-11.c new file mode 100644 index 00000000000..82ed09d9314 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-11.c @@ -0,0 +1,73 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ +/* { dg-xfail-run-if "" { offload_device_nonshared_as } } */ + +/* Two levels of pointer indirection dereferenced inside a target region cause + an illegal memory access, even without noncontiguous updates. */ + +#include +#include + +#define DIM1 4 +#define DIM2 3 +#define N 20 + +int +main () +{ + int **arr[DIM1]; + int i, j, k; + + for (i = 0; i < DIM1; i++) + { + arr[i] = (int **) malloc (DIM2 * sizeof (int *)); + for (j = 0; j < DIM2; j++) + arr[i][j] = (int *) calloc (N, sizeof (int)); + } + + /* Manual deep copy. */ +#pragma omp target enter data map(alloc: arr[0:DIM1]) + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(alloc: arr[i][0:DIM2]) + } + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target enter data map(to: arr[i][j][0:N]) + } + + /* Crash here: dereferencing two levels of indirection to write + data from target code. */ +#pragma omp target map(alloc: arr) + { + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + for (int k = 0; k < N; k++) + arr[i][j][k] = i * 10000 + j * 100 + k; + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target exit data map(release: arr[i][j]) map(from: arr[i][j][0:N]) + } + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(release: arr[i], arr[i][0:DIM2]) + } +#pragma omp target exit data map(release: arr, arr[0:DIM1]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (k = 0; k < N; k++) + assert (arr[i][j][k] == i * 10000 + j * 100 + k); + + for (i = 0; i < DIM1; i++) + { + for (j = 0; j < DIM2; j++) + free (arr[i][j]); + free (arr[i]); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-2.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-2.c new file mode 100644 index 00000000000..a8982133a95 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-2.c @@ -0,0 +1,75 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Noncontiguous "target update" through a 2D array of pointers where the + pointers themselves are the mapped payload rather than a point of indirection + -- one segment, two dimensions, with the final pointer left un-dereferenced. + */ + +#include + +#define DIM1 5 +#define DIM2 5 + +int main() +{ + int *y[DIM1][DIM2]; + int markers[2]; + int i, j; + + /* Case 1: update to device, a partial range of one fixed row. */ + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + y[i][j] = 0; + +#pragma omp target enter data map(to: y) + + for (j = 0; j < 2; j++) + y[0][j] = &markers[j]; + +#pragma omp target update to(y[0][:2]) + + /* Erase the host copy before pulling back from the device, so the + assertions below only see what actually made it across. */ + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + y[i][j] = 0; + +#pragma omp target exit data map(from: y) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i == 0 && j < 2) + assert (y[i][j] == &markers[j]); + else + assert (y[i][j] == 0); + + /* Case 2: update from device, same shape, on a different row. */ + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + y[i][j] = 0; + +#pragma omp target enter data map(to: y) + + for (j = 0; j < 2; j++) + y[3][j] = &markers[j]; + +#pragma omp target update to(y[3][:2]) + + for (j = 0; j < 2; j++) + y[3][j] = 0; + +#pragma omp target update from(y[3][:2]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + if (i == 3 && j < 2) + assert (y[i][j] == &markers[j]); + else + assert (y[i][j] == 0); + +#pragma omp target exit data map(delete: y) + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-3.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-3.c new file mode 100644 index 00000000000..10eb13d600a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-3.c @@ -0,0 +1,135 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test noncontiguous "target update" through one level of array-shaped pointer + indirection, with two dimensions on the array-of-pointers side. + Also check the transferred data can be accessed from device code. */ + +#include +#include +#include + +#define DIM1 3 +#define DIM2 4 +#define ROWS 6 +#define COLS 6 + +int main() +{ + int *x[DIM1][DIM2]; + int i, j, r, c; + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + x[i][j] = (int *) calloc (ROWS * COLS, sizeof (int)); + +#pragma omp target enter data map(alloc: x[0:DIM1]) + + /* Case 1: update to device. Two fixed indices select a single pointer + from the 2D array-of-pointers (segment 0); the selected pointer + contributes a 2D sub-rectangle of its buffer (segment 1's two + dimensions), the row dimension selected with a stride of 2 (rows + 1, 3 and 5). */ + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target enter data map(to: x[i][j][:ROWS*COLS]) + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x[1][1][r * COLS + c] = 100 * r + c; + +#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2][0:2]) + +#pragma omp target map(alloc: x) + { + for (int r = 1; r <= 5; r += 2) + for (int c = 0; c < 2; c++) + x[1][1][r * COLS + c] += 1000; + } + + /* Erase the host copies before pulling back from the device, so the + assertions below only see what actually made it across. */ + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + memset (x[i][j], 0, ROWS * COLS * sizeof (int)); + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target exit data map(from: x[i][j][:ROWS*COLS]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + { + int expected = 0; + if (i == 1 && j == 1 && (r == 1 || r == 3 || r == 5) && c < 2) + expected = 100 * r + c + 1000; + assert (x[i][j][r * COLS + c] == expected); + } + + /* Case 2: update from device, same two-segment/two-dimension shape, on + a different pointer: clobber the host copy after pushing data to the + device, then confirm "update from" actually pulls the device's data + back rather than leaving the clobbered value. */ + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + memset (x[i][j], 0, ROWS * COLS * sizeof (int)); + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target enter data map(to: x[i][j][:ROWS*COLS]) + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x[2][0][r * COLS + c] = 300 + 10 * r + c; + +#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2][1:3]) + +#pragma omp target map(alloc: x) + { + for (int r = 2; r <= 3; r++) + for (int c = 1; c <= 3; c++) + x[2][0][r * COLS + c] += 1000; + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x[2][0][r * COLS + c] = -1; + +#pragma omp target update from((([ROWS][COLS]) x[2][0:1])[2:2][1:3]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + { + if (i == 2 && j == 0 && r >= 2 && r <= 3 && c >= 1 && c <= 3) + assert (x[i][j][r * COLS + c] == 300 + 10 * r + c + 1000); + else if (i == 2 && j == 0) + assert (x[i][j][r * COLS + c] == -1); + else + assert (x[i][j][r * COLS + c] == 0); + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target exit data map(delete: x[i][j][:ROWS*COLS]) + } + +#pragma omp target exit data map(release: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + free (x[i][j]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-4.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-4.c new file mode 100644 index 00000000000..39afb53e216 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-4.c @@ -0,0 +1,134 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test noncontiguous 'target update" through one level of pointer indirection, + where the array-of-pointers side has several dimensions selected by fixed + indices. Also check the transferred data can be accessed from device + code. */ + +#include +#include +#include + +#define D1 3 +#define D2 3 +#define D3 4 +#define M 8 + +int main() +{ + int *w[D1][D2][D3]; + int i, j, k, m; + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + w[i][j][k] = (int *) calloc (M, sizeof (int)); + +#pragma omp target enter data map(alloc: w[0:D1]) + + /* Case 1: update to device. Three consecutive fixed indices select one + specific pointer, then a partial range of its buffer is updated. */ + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + { +#pragma omp target enter data map(to: w[i][j][k][:M]) + } + + for (m = 0; m < M; m++) + w[1][2][0][m] = 10 + m; + +#pragma omp target update to(w[1][2][0][2:4]) + +#pragma omp target map(alloc: w) + { + for (int m = 2; m < 6; m++) + w[1][2][0][m] += 1000; + } + + /* Erase the host copies before pulling back from the device, so the + assertions below only see what actually made it across. */ + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + memset (w[i][j][k], 0, M * sizeof (int)); + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + { +#pragma omp target exit data map(from: w[i][j][k][:M]) + } + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + for (m = 0; m < M; m++) + if (i == 1 && j == 2 && k == 0 && m >= 2 && m < 6) + assert (w[i][j][k][m] == 10 + m + 1000); + else + assert (w[i][j][k][m] == 0); + + /* Case 2: update from device, same shape (three consecutive fixed indices), + through a different pointer: clobber the host copy after pushing data + to the device, then confirm "update from" actually pulls the + device's data back rather than leaving the clobbered value. */ + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + memset (w[i][j][k], 0, M * sizeof (int)); + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + { +#pragma omp target enter data map(to: w[i][j][k][:M]) + } + + for (m = 0; m < M; m++) + w[2][0][3][m] = 100 + m; + +#pragma omp target update to(w[2][0][3][1:5]) + +#pragma omp target map(alloc: w) + { + for (int m = 1; m < 6; m++) + w[2][0][3][m] += 1000; + } + + for (m = 0; m < M; m++) + w[2][0][3][m] = -1; + +#pragma omp target update from(w[2][0][3][1:5]) + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + for (m = 0; m < M; m++) + { + if (i == 2 && j == 0 && k == 3 && m >= 1 && m < 6) + assert (w[i][j][k][m] == 100 + m + 1000); + else if (i == 2 && j == 0 && k == 3) + assert (w[i][j][k][m] == -1); + else + assert (w[i][j][k][m] == 0); + } + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + { +#pragma omp target exit data map(delete: w[i][j][k][:M]) + } + +#pragma omp target exit data map(release: w[0:D1]) + + for (i = 0; i < D1; i++) + for (j = 0; j < D2; j++) + for (k = 0; k < D3; k++) + free (w[i][j][k]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-5.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-5.c new file mode 100644 index 00000000000..63d798aaa7c --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-5.c @@ -0,0 +1,136 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test noncontiguous target update through one level of pointer indirection, + where the array-shaping cast has more dimensions than are explicitly + sectioned afterwards. Also check that transferred data can be read and + written from device code. */ + +#include +#include +#include + +#define DIM1 3 +#define DIM2 4 +#define ROWS 6 +#define COLS 6 + +int main() +{ + int *x[DIM1][DIM2]; + int i, j, r, c; + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + x[i][j] = (int *) calloc (ROWS * COLS, sizeof (int)); + +#pragma omp target enter data map(alloc: x[0:DIM1]) + + /* Case 1: update to device. Two fixed indices select a single pointer + from the 2D array-of-pointers (segment 0); the selected pointer + contributes a strided range of whole rows (segment 1's only explicit + dimension). */ + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target enter data map(to: x[i][j][:ROWS*COLS]) + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x[1][1][r * COLS + c] = 100 * r + c; + +#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2]) + +#pragma omp target map(alloc: x) + { + for (int r = 1; r <= 5; r += 2) + for (int c = 0; c < COLS; c++) + x[1][1][r * COLS + c] += 1000; + } + + /* Erase the host copies before pulling back from the device, so the + assertions below only see what actually made it across. */ + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + memset (x[i][j], 0, ROWS * COLS * sizeof (int)); + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target exit data map(from: x[i][j][:ROWS*COLS]) + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + { + int expected = 0; + if (i == 1 && j == 1 && (r == 1 || r == 3 || r == 5)) + expected = 100 * r + c + 1000; + assert (x[i][j][r * COLS + c] == expected); + } + + /* Case 2: update from device, same shape (rows selected, columns left + to their whole span), on a different pointer: clobber the host copy + after pushing data to the device, then confirm "update from" actually + pulls the device's data back rather than leaving the clobbered + value. */ + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + memset (x[i][j], 0, ROWS * COLS * sizeof (int)); + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target enter data map(to: x[i][j][:ROWS*COLS]) + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x[2][0][r * COLS + c] = 300 + 10 * r + c; + +#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2]) + +#pragma omp target map(alloc: x) + { + for (int r = 2; r <= 3; r++) + for (int c = 0; c < COLS; c++) + x[2][0][r * COLS + c] += 1000; + } + + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + x[2][0][r * COLS + c] = -1; + +#pragma omp target update from((([ROWS][COLS]) x[2][0:1])[2:2]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (r = 0; r < ROWS; r++) + for (c = 0; c < COLS; c++) + { + if (i == 2 && j == 0 && r >= 2 && r <= 3) + assert (x[i][j][r * COLS + c] == 300 + 10 * r + c + 1000); + else if (i == 2 && j == 0) + assert (x[i][j][r * COLS + c] == -1); + else + assert (x[i][j][r * COLS + c] == 0); + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target exit data map(delete: x[i][j][:ROWS*COLS]) + } + +#pragma omp target exit data map(release: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + free (x[i][j]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-6.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-6.c new file mode 100644 index 00000000000..01127006d55 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-6.c @@ -0,0 +1,75 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test array-shaping cast applied to a plain pointer, with a further array + section applied on top of the cast. */ + +#include +#include +#include + +#define N 100 + +int main() +{ + int *ptr = (int *) calloc (N, sizeof (int)); + int i; + +#pragma omp target enter data map(to: ptr[:N]) + + for (i = 0; i < N; i++) + ptr[i] = i; + +#pragma omp target update to((([N]) ptr)[10:30]) + +#pragma omp target map(alloc: ptr[:N]) + { + for (int i = 10; i < 40; i++) + ptr[i] += 1000; + } + + /* Erase the host copy before pulling back from the device, so the + assertions below only see what actually made it across. */ + memset (ptr, 0, N * sizeof (int)); + +#pragma omp target exit data map(from: ptr[:N]) + + for (i = 0; i < N; i++) + if (i >= 10 && i < 40) + assert (ptr[i] == i + 1000); + else + assert (ptr[i] == 0); + + /* Update from device: clobber the host copy after pushing data to the + device, then confirm "update from" actually pulls the device's data + back rather than leaving the clobbered value. */ + +#pragma omp target enter data map(to: ptr[:N]) + + for (i = 0; i < N; i++) + ptr[i] = i; + +#pragma omp target update to((([N]) ptr)[10:30]) + +#pragma omp target map(alloc: ptr[:N]) + { + for (int i = 10; i < 40; i++) + ptr[i] += 1000; + } + + for (i = 0; i < N; i++) + ptr[i] = -1; + +#pragma omp target update from((([N]) ptr)[10:30]) + + for (i = 0; i < N; i++) + if (i >= 10 && i < 40) + assert (ptr[i] == i + 1000); + else + assert (ptr[i] == -1); + +#pragma omp target exit data map(delete: ptr[:N]) + + free (ptr); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-7.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-7.c new file mode 100644 index 00000000000..b75ccea2539 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-7.c @@ -0,0 +1,104 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test array-shaping cast applied directly to a fixed-index element of a 1D + array-of-pointers, with no further outer section.. */ + +#include +#include +#include + +#define DIM1 5 +#define N 20 + +int main () +{ + int *x[DIM1]; + int i, k; + + for (i = 0; i < DIM1; i++) + x[i] = (int *) calloc (N, sizeof (int)); + +#pragma omp target enter data map(alloc: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:N]) + } + + /* Case 1: update to device, single fixed pointer index, whole span. */ + + for (k = 0; k < N; k++) + x[2][k] = 100 + k; + +#pragma omp target update to(([N]) x[2]) + +#pragma omp target map(alloc: x) + { + for (int k = 0; k < N; k++) + x[2][k] += 1000; + } + + for (i = 0; i < DIM1; i++) + memset (x[i], 0, N * sizeof (int)); + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(from: x[i][:N]) + } + + for (i = 0; i < DIM1; i++) + for (k = 0; k < N; k++) + { + int expected = (i == 2) ? 100 + k + 1000 : 0; + assert (x[i][k] == expected); + } + + /* Case 2: update from device, same shape, different index, confirming + "update from" actually pulls the device's data back. */ + + for (i = 0; i < DIM1; i++) + { +#pragma omp target enter data map(to: x[i][:N]) + } + + for (k = 0; k < N; k++) + x[3][k] = 200 + k; + +#pragma omp target update to(([N]) x[3]) + +#pragma omp target map(alloc: x) + { + for (int k = 0; k < N; k++) + x[3][k] += 1000; + } + + for (k = 0; k < N; k++) + x[3][k] = -1; + +#pragma omp target update from(([N]) x[3]) + + for (i = 0; i < DIM1; i++) + for (k = 0; k < N; k++) + { + int expected; + if (i == 3) + expected = 200 + k + 1000; + else if (i == 2) + expected = 100 + k + 1000; + else + expected = 0; + assert (x[i][k] == expected); + } + + for (i = 0; i < DIM1; i++) + { +#pragma omp target exit data map(delete: x[i][:N]) + } + +#pragma omp target exit data map(release: x[0:DIM1]) + + for (i = 0; i < DIM1; i++) + free (x[i]); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-8.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-8.c new file mode 100644 index 00000000000..eccc187dbd6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-8.c @@ -0,0 +1,87 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test three-segment (two pointer indirections) noncontiguous target update: + a strided section reached by crossing two levels of pointer indirection. + Only one dimension per segment. */ + +#include +#include + +#define DIM1 4 +#define DIM2 3 +#define N 20 + +int +main () +{ + int **arr[DIM1]; + int i, j, k; + + for (i = 0; i < DIM1; i++) + { + arr[i] = (int **) malloc (DIM2 * sizeof (int *)); + for (j = 0; j < DIM2; j++) + arr[i][j] = (int *) calloc (N, sizeof (int)); + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (k = 0; k < N; k++) + arr[i][j][k] = i * 10000 + j * 100 + k; + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target enter data map(to: arr[i][j][:N]) + } + + /* Corrupt the host copy of exactly the two rows the update below will + touch, so matching the original values afterward proves the update + actually pulled from the device instead of being a no-op. */ + for (j = 0; j < DIM2; j += 2) + for (k = 0; k < N; k++) + arr[2][j][k] = -1; + +#pragma omp target update from(arr[2][0:2:2][3:5:2]) + + { + int touched[5] = { 3, 5, 7, 9, 11 }; + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + for (k = 0; k < N; k++) + { + int expected; + if (i == 2 && (j == 0 || j == 2)) + { + int is_touched = 0, t; + for (t = 0; t < 5; t++) + if (touched[t] == k) + is_touched = 1; + /* Restored from the device if touched, still corrupted + otherwise -- proves the update copied exactly the + requested elements. */ + expected = is_touched ? i * 10000 + j * 100 + k : -1; + } + else + /* Untouched by either the corruption or the update. */ + expected = i * 10000 + j * 100 + k; + assert (arr[i][j][k] == expected); + } + } + + for (i = 0; i < DIM1; i++) + for (j = 0; j < DIM2; j++) + { +#pragma omp target exit data map(delete: arr[i][j][:N]) + } + + for (i = 0; i < DIM1; i++) + { + for (j = 0; j < DIM2; j++) + free (arr[i][j]); + free (arr[i]); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-section-9.c b/libgomp/testsuite/libgomp.c-c++-common/array-section-9.c new file mode 100644 index 00000000000..6fc152dcd1f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/array-section-9.c @@ -0,0 +1,120 @@ +/* { dg-do run { target offload_device_nonshared_as } } */ + +/* Test three-segment (two pointer indirections) noncontiguous target update, + with *two* dimensions selected per segment -- six dimensions total. */ + +#include +#include + +#define DIM1A 2 +#define DIM1B 3 +#define ROWS1 3 +#define NCOLS 3 +#define ROWS2 3 +#define ROWLEN 20 + +typedef int (*p2_t)[ROWLEN]; +typedef p2_t (*p1_t)[NCOLS]; + +static int +encode (int i, int j, int r1, int c1, int r2, int k) +{ + return i * 1000000 + j * 100000 + r1 * 10000 + c1 * 1000 + r2 * 100 + k; +} + +int +main () +{ + p1_t arr[DIM1A][DIM1B]; + int i, j, r1, c1, r2, k; + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + { + arr[i][j] = (p1_t) malloc (ROWS1 * NCOLS * sizeof (p2_t)); + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + arr[i][j][r1][c1] + = (p2_t) malloc (ROWS2 * ROWLEN * sizeof (int)); + } + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + for (r2 = 0; r2 < ROWS2; r2++) + for (k = 0; k < ROWLEN; k++) + arr[i][j][r1][c1][r2][k] = encode (i, j, r1, c1, r2, k); + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + { +#pragma omp target enter data map(to: arr[i][j][r1][c1][:ROWS2]) + } + + /* Corrupt the whole of ARR[1]'s data (every J/R1/C1/R2/K), so that + precisely matching the six-dimensional touched set below after the + update proves each dimension of each segment was walked correctly. */ + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + for (r2 = 0; r2 < ROWS2; r2++) + for (k = 0; k < ROWLEN; k++) + arr[1][j][r1][c1][r2][k] = -1; + +#pragma omp target update from(arr[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2]) + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + for (r2 = 0; r2 < ROWS2; r2++) + for (k = 0; k < ROWLEN; k++) + { + int expected; + if (i == 1) + { + int j_touched = (j == 0 || j == 2); + int r1_touched = (r1 == 0 || r1 == 2); + int c1_touched = (c1 == 0 || c1 == 2); + int r2_touched = (r2 == 0 || r2 == 1); + int k_touched + = (k == 3 || k == 5 || k == 7 || k == 9 || k == 11); + if (j_touched && r1_touched && c1_touched + && r2_touched && k_touched) + /* Restored from the device -- proves this exact + combination of dimensions across all three + segments was selected. */ + expected = encode (i, j, r1, c1, r2, k); + else + /* Still corrupted -- proves the update did not + touch anything outside the requested set. */ + expected = -1; + } + else + /* Untouched by either the corruption or the update. */ + expected = encode (i, j, r1, c1, r2, k); + assert (arr[i][j][r1][c1][r2][k] == expected); + } + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + { +#pragma omp target exit data map(delete: arr[i][j][r1][c1][:ROWS2]) + } + + for (i = 0; i < DIM1A; i++) + for (j = 0; j < DIM1B; j++) + { + for (r1 = 0; r1 < ROWS1; r1++) + for (c1 = 0; c1 < NCOLS; c1++) + free (arr[i][j][r1][c1]); + free (arr[i][j]); + } + + return 0; +}