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; +}