From patchwork Fri Jul 26 18:05:43 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1965380 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=fail reason="signature verification failed" (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=PQAGur4Y; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4WVwg06fjcz1yXx for ; Sat, 27 Jul 2024 04:06:20 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 1CB503858CD9 for ; Fri, 26 Jul 2024 18:06:18 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wr1-x429.google.com (mail-wr1-x429.google.com [IPv6:2a00:1450:4864:20::429]) by sourceware.org (Postfix) with ESMTPS id 5968C3858CD9 for ; Fri, 26 Jul 2024 18:05:47 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 5968C3858CD9 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 5968C3858CD9 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::429 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1722017150; cv=none; b=ZHwnAdyBxfmxMmPPNrlxZOgU+1QQmCsTzkQ1dYJPQHhG+ezu4yollICTvt3w07K04oqhWhQKzdtWG+VbykbB6tszY+JSAjxzDvnHvEJmaXwz5G4VMruYmzM2A/fFrkZaZOrKLnqDjBt8Jns+KJZbtUsasfjBCi4dj0AmpL7iOuY= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1722017150; c=relaxed/simple; bh=FtAvfJPO2uZ9XvcShqwWK29IfnAPuTN8tLy4xJONdaU=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:From:Subject; b=hNWczUEZMw+0NY7JPvO+zSU0zy4aoNJg69sqc/K3yIJ4AOb1Q3CUUItsN0089Qr+JO2Uih7H1GPGln3owKA53JEDUnJecvmlZOY0VyZzYXZYVQLH7BBu75lW0XBzcq7fNyuA4juE7bfjj2mx1Q65LPrRS4iiJ+MNH6e8aEzN1c0= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wr1-x429.google.com with SMTP id ffacd0b85a97d-3686b285969so1457134f8f.0 for ; Fri, 26 Jul 2024 11:05:47 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1722017146; x=1722621946; darn=gcc.gnu.org; h=subject:from:to:content-language:user-agent:mime-version:date :message-id:from:to:cc:subject:date:message-id:reply-to; bh=Durv/2Pymssq16oiE/+ZYJkOZ/fxqulybZ0c8LfTx94=; b=PQAGur4Y0s5mTdTrEKdgjfcQ3jd3GNWk+jIbAFNI8Q1ZNmC1FH6qKUAogbs1t/a9ZY n2FaUiW1a7nOo7KYEcpJEPayqe6e3xVA0rfgWzvEO3GmTKClYdgWV6FTZO464HZjsBNm 3GEZrJhvj7jP9yxv4LMPi9N8GKAPW2vJyralvs9pCJn7n4kDdOtDKdX9iYOVofOmc6TZ MAHUgGflaeAg/vzuDvta4TAqEEkp8Rk5QRhUN5Y3/UVvQGuz80QEpxp3mLdd7UXcypau 3GLonVX/c/GD/o251dtubuJ0fH6gI6LuD/8MTLeNVzRgU8k4I4l5tmVy2pa2viZEKevk D6bg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1722017146; x=1722621946; h=subject:from:to:content-language:user-agent:mime-version:date :message-id:x-gm-message-state:from:to:cc:subject:date:message-id :reply-to; bh=Durv/2Pymssq16oiE/+ZYJkOZ/fxqulybZ0c8LfTx94=; b=tpg/DkkitGTuxyGfkNUk02y9H9bVh3Eyh6mmHUvNMJrf4JgWkCbZ2cep29fGZIW+DB KvgSMT6o3HzpKgAxX5uuApbsRqDosdHzDw6bLEE9DDer/Nd6BAyeJ1FYv2naWSHDi3NP T1WVLzGFiBpfOOAb+zlFtnEfogoanc3PBrbVmVBA257KJY8YW2q4kMm9wOa2/FM3zrOF yUPFI6BRUdmtmfC01RZpUDrEZlcG+ub4cxgkwQpA9D1mo9QhWl6KGSGgr9c3U2joF1Mq x57moRp35+2y3oPvdKUg+JnbAJBspa++SevStBwqGxvG+LFyX7CbX+O+Tf/HYiLXBo1B fEpA== X-Gm-Message-State: AOJu0YxUFQvmOpUTU1Zuv1EO5thFD9nThspVqsXdEhbVoop8i7zzznsq rdynAk5QLO7smtVkQ8Ohx1MXvl8qyRvhJ2iLeYlqqbo96FGg5Kwg/BsvIjGtP5fNAIr9rkV0AA9 U+qY= X-Google-Smtp-Source: AGHT+IGg8777atLEmuw4i3RGJaP8/lH8G9TpvyQWI2X/nj4wdgovrw9vjti+OY0BrSAJqLJunsgbuw== X-Received: by 2002:a05:6000:1a43:b0:367:9d82:8370 with SMTP id ffacd0b85a97d-36b5d073fa6mr285078f8f.45.1722017145695; Fri, 26 Jul 2024 11:05:45 -0700 (PDT) Received: from [192.168.8.100] (tmo-085-80.customers.d1-online.com. [80.187.85.80]) by smtp.gmail.com with ESMTPSA id ffacd0b85a97d-36b36801828sm5698217f8f.65.2024.07.26.11.05.44 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 26 Jul 2024 11:05:45 -0700 (PDT) Message-ID: <9b132853-8fce-4433-9683-c2cba31367c2@baylibre.com> Date: Fri, 26 Jul 2024 20:05:43 +0200 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Content-Language: en-US To: gcc-patches , Jakub Jelinek From: Tobias Burnus Subject: [Patch] libgomp: Fix declare target link with offset array-section mapping [PR116107] X-Spam-Status: No, score=-11.2 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, HTML_MESSAGE, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.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~incoming=patchwork.ozlabs.org@gcc.gnu.org The main idea of 'link' is to permit putting only a subset of a huge array on the device. Well, in order to make this work properly, it requires that one can map an array section, which does not start with the first element. This patch adjusts the pointers such, that this actually works. (Tested on x86-64-gnu-linux with Nvptx offloading.) Comments, suggestions, remarks before I commit it? Tobias libgomp: Fix declare target link with offset array-section mapping [PR116107] Assume that 'int var[100]' is 'omp declare target link(var)'. When now mapping an array section with offset such as 'map(to:var[20:10])', the device-side link pointer has to store &[0] minus the offset such that var[20] will access [0]. But the offset calculation was missed such that the device-side 'var' pointed to the first element of the mapped data - and var[20] points beyond at some invalid memory. PR middle-end/116107 libgomp/ChangeLog: * target.c (gomp_map_vars_internal): Honor array mapping offsets with declare-target 'link' variables. * testsuite/libgomp.c-c++-common/target-link-2.c: New test. libgomp/target.c | 7 ++- .../testsuite/libgomp.c-c++-common/target-link-2.c | 59 ++++++++++++++++++++++ 2 files changed, 64 insertions(+), 2 deletions(-) diff --git a/libgomp/target.c b/libgomp/target.c index aa01c1367b9..e3e648f5443 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1820,8 +1820,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, if (k->aux && k->aux->link_key) { /* Set link pointer on target to the device address of the - mapped object. */ - void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset); + mapped object. Also deal with offsets due to + array-section mapping. */ + void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset + - (k->host_start + - k->aux->link_key->host_start)); /* We intentionally do not use coalescing here, as it's not data allocated by the current call to this function. */ gomp_copy_host2dev (devicep, aq, (void *) n->tgt_offset, diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c new file mode 100644 index 00000000000..4ff4080da76 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c @@ -0,0 +1,59 @@ +/* PR middle-end/116107 */ + +#include + +int arr[15] = {1,2,3,4,5,6,7,8,9,10,11,12,13,14,15}; +#pragma omp declare target link(arr) + +#pragma omp begin declare target +void f(int *res) +{ + __builtin_memcpy (res, &arr[5], sizeof(int)*10); +} + +void g(int *res) +{ + __builtin_memcpy (res, &arr[3], sizeof(int)*10); +} +#pragma omp end declare target + +int main() +{ + int res[10], res2; + for (int dev = 0; dev < omp_get_num_devices(); dev++) + { + __builtin_memset (res, 0, sizeof (res)); + res2 = 99; + + #pragma omp target enter data map(arr[5:10]) device(dev) + + #pragma omp target map(from: res) device(dev) + f (res); + + #pragma omp target map(from: res2) device(dev) + res2 = arr[5]; + + if (res2 != 6) + __builtin_abort (); + for (int i = 0; i < 10; i++) + if (res[i] != 6 + i) + __builtin_abort (); + + #pragma omp target exit data map(release:arr[5:10]) device(dev) + + for (int i = 0; i < 15; i++) + res[i] *= 10; + __builtin_abort (); + + #pragma omp target enter data map(arr[3:10]) device(dev) + __builtin_memset (res, 0, sizeof (res)); + + #pragma omp target map(from: res) device(dev) + g (res); + + for (int i = 0; i < 10; i++) + if (res[i] != (4 + i)*10) + __builtin_abort (); + } + return 0; +}