From patchwork Wed Sep 2 12:07:56 2026 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 142705 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 F20864BA2E21 for ; Wed, 2 Sep 2026 12:09:34 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org F20864BA2E21 Authentication-Results: sourceware.org; dkim=fail reason="signature verification failed" (2048-bit key, secure) header.d=baylibre.com header.i=@baylibre.com header.a=rsa-sha256 header.s=google header.b=nK4uVqzm X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x335.google.com (mail-wm1-x335.google.com [IPv6:2a00:1450:4864:20::335]) by sourceware.org (Postfix) with ESMTPS id A748A4BA23E4 for ; Wed, 2 Sep 2026 12:08:00 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org A748A4BA23E4 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 A748A4BA23E4 Authentication-Results: sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::335 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1788350894; cv=none; b=S104GOYON7LgzEk5lgHySJiW9YNQI6cOjDa8ehPwlPCnQYR5hhMjO5ueb0cCeH6EKktoRnmxJ85mRnIt8MiyCQDBS3Ho8d419LR8h5aUwyWtK5f3Cqat3MvUeJ4k348gVZ4jNXCVmyttyVEwQKkPYgLAl0kWUul8FqUhTvGqSKw= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1788350894; c=relaxed/simple; bh=FhQ7/ca5oFupOVS2QzbMwPw4t48+X4UE2UzxRe9dI2U=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:From:Subject; b=tAe5QKkddBAHtxZz+2OUBc6nJMDyZDegn2AaKFogjPYxMfrTwwZuD45kLTZWwtAUXc1jaPUP8u6xN9Fd7ZWmJtcwGXpY9XOZoPIWOyoKATu1FWtkbHixmRKih6jncRgmscLuBfwX/o1rKv/5CwKXFclZvpxDRZARtzca/5x10Lk= 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=nK4uVqzm DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A748A4BA23E4 Received: by mail-wm1-x335.google.com with SMTP id 5b1f17b1804b1-495437bb891so7811485e9.1 for ; Wed, 02 Sep 2026 05:08:00 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre.com; s=google; t=1788350879; x=1788955679; darn=gcc.gnu.org; h=subject:from:to:content-language:user-agent:mime-version:date :message-id:content-type:from:to:cc:subject:date:message-id:reply-to :content-type; bh=0yiMQnv4DrJVwzXtb0VhS6ExPxWYMYHO5HMAHB3tjMk=; b=nK4uVqzmNEN3d3H84Yrlfj9EnAf9xwIm1STFf+CWbaF3ME9N2O4tNfduHsgUFGS6Jm vxxAih+WasvE7nly8s+UhitHLwb00rYD69B+zt6jYqUehqco3yPAU9eV7XLMlKuSGA+2 lM24zqrOXpDADcpXizwGcjp0ilQK36GKDQyOruzcRGt4LtCctUPKSsM4kqYM21Lvmm/b WPybbT32Cm/QOeeydSqYE5o0FFhXOnSgLuhmFUBksOHPD1Sor5FbEPh+jl5njcwA6vy0 Wq9PTADGZyhnIv3gHvTXUkulA0suDC1sCZha7HGDJ1HKxBIj/KtCC8NJ4uKr0ut3TiMd rAcQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20251104; t=1788350879; x=1788955679; h=subject:from:to:content-language:user-agent:mime-version:date :message-id:content-type:x-gm-gg:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to:content-type; bh=0yiMQnv4DrJVwzXtb0VhS6ExPxWYMYHO5HMAHB3tjMk=; b=rL8zu6fC1LorXX1N6Z97pukW1/fQDd4R5IzQQw+K14hzpa+oMGkuxAPfkqLeJUQG0q w3jj1TH3iDurrpWpQfpC+Vfk/tsMMqFBWAkaVve/gwC9OVcofaQdgYBMEQJt3zh0sTIy MTNDJfXBk/g4XL3K0rqPuNnehr/TEyC4eaNZZF2gGT3u3EgaLQo6RfSQUfWrErnlH0RK Ud2vA8/XDW6ZK6Fs4SCHa8+YSJxACH13dCzu3cMfTJDPKi/0TK+N12gXcKWUWkmrNA4i k00MF5+ffNsLEUNHyqZKv2poaMVCtJvH/Cpl1wMBJjBQUxNJvlGn17/1ucTOYLal265h BJiw== X-Gm-Message-State: AFuF++nabfVboCCBCbO3laMhz/oeviVxm9hDp+7HDK1r2DIVkoyHOH8G dZ+q8TY2mHdv3QfPqQSGky1EWHmfAqfgycWCYLQDs1op3+Mdua853eRsgw7TZxgoKaFiZmps+XK hv3FM X-Gm-Gg: AR+sD13s2dxDt0NhiEI2FaPWZ5C47yzlCrj71XcJisTt3DOyvm+MUIrkcCDA+BX3gJS X8YUd7v/FsLz4z+o8mA5RgItNILozzLkOqM5Gsc8xAzGtalukXe+1/K++Kzuwe4J8a089ccoiau gqt+hfrP30zvQ1wOmeazgu0h/1+6E+U2/fo7l7ADqavua9/73mnhzgxouxWEPjr6TF5R1JtSt87 oy7hk84PUcFNGjLfUjgDSuAxtuS3fSXFZd/036OWjU0CgoEp6rO0wH9Z26DkJRZo6qvO42ckW6h iAvJiXT2X2rCL6EJsOnFqypsPKw6ompQiR/YgBiXWaTL1r7AMI6pb2Nds51OuhatFay+PrquLV+ NxrcWtHmHyRQYMi0rQ92WvpLOI8OuLhSAxrSkwXKH3CFCXXPr5/GrLpzIULojphfl1clUFqHH/3 dKGGiKrE6GbQr0DTs6F/e/JtXgIGd3pNW2N/fAUqHeadiqLZUFsff4XeNpcg4wfI4uzVCmUOb6d JcYdhP2pKSoUn/7j/xQeYejxUSLjSqVQuGP+iknMMA2gEYAksfatlkKf9Lp03U0iF5PiYtOVgbx CMYZNPX7zZjiwwF+iShpXls= X-Received: by 2002:a05:600c:8b6d:b0:49c:e2b3:7474 with SMTP id 5b1f17b1804b1-49ce7b66b54mr44326715e9.0.1788350878370; Wed, 02 Sep 2026 05:07:58 -0700 (PDT) Received: from ?IPV6:2001:16b8:2603:8b00:f59e:cf4f:df4f:9a18? (200116b826038b00f59ecf4fdf4f9a18.dip.versatel-1u1.de. [2001:16b8:2603:8b00:f59e:cf4f:df4f:9a18]) by smtp.gmail.com with ESMTPSA id 5b1f17b1804b1-49ce7d04c69sm31736225e9.8.2026.09.02.05.07.57 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Wed, 02 Sep 2026 05:07:57 -0700 (PDT) Message-ID: <8f9480fa-9ddd-49a7-9b6b-1034bdbdc74f@baylibre.com> Date: Wed, 2 Sep 2026 14:07:56 +0200 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Content-Language: en-US To: gcc-patches From: Tobias Burnus Subject: [committed] libgomp: fix comment, add testcase libgomp.c/usm_env_handling-2.c [PR127186] X-Spam-Status: No, score=-11.6 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, HTML_MESSAGE, RCVD_IN_DNSWL_NONE, RCVD_IN_PBL, 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 There are some known issues with GOMP_RUNTIME_USM=enabled where the a global static device variable is not updated, which leads to wrong results. Commit r17-3807-g82104cde7ac0ba added a check for this, such that when those variables appear with the 'always' modifier an error was printed. I wrongly claimed in the patch email and comments (and the here referenced PR) that for 'target map(var)' it will just pass through the host address, which causes consistency issues. However, that turned out to be not the case as GCC directly uses the variable. This commit fixes the comment and adds a testcase for it. r17-3858-gb590dc3fdc0aeb Note: The issue with 'always' still exists (main reason for PR127186), i.e. that the code that is conditionally disabled in libgomp.c/usm_env_handling-2.c still fails because of this bug / missing feature. See PR for details. Tobias commit b590dc3fdc0aeb9bcd923cfd22778a0864102101 Author: Tobias Burnus Date: Wed Sep 2 13:57:30 2026 +0200 libgomp: fix comment, add testcase libgomp.c/usm_env_handling-2.c [PR127186] The following is with GOMP_RUNTIME_USM=enabled and a suitable device available: In r17-3807-g82104cde7ac0ba I wrongly claimed that in a 'target' structured block, a global static variable will not be address translated such that the code will access the host variable, which would cause inconsistency issues (aka wrong code). However, the compiler middle end already optimizes his by accessing directly the respective global variable and not even passing the variable address to libgomp for host->device address translation. Thus, there is no issue in this regard. Add a testcase for this. libgomp/ChangeLog: PR libgomp/127186 * target.c (GOMP_target_ext): Fix comment related to GOMP_RUNTIME_USM. * testsuite/libgomp.c/usm_env_handling-2.c: New test. --- libgomp/target.c | 6 +- libgomp/testsuite/libgomp.c/usm_env_handling-2.c | 157 +++++++++++++++++++++++ 2 files changed, 160 insertions(+), 3 deletions(-) diff --git a/libgomp/target.c b/libgomp/target.c index f688f33378d..081701ca634 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -3719,9 +3719,9 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, need to be handled in map clauses. TODO: Implement this feature; for now fail with an error if those are encountered. - FIXME: Not only 'always' has to be handled as variables in the target - region and an (indirect) access to the global variable in a function - call must refer to the same variable. */ + Note that only the data transfer it required. Code inside target + structured blocks already directly accesses device variable; + 'hostaddrs' does not even contain those for mapping purposes. */ constexpr bool short_mapkind = true; /* OpenMP */ constexpr int typemask = 0xff; for (size_t i = 0; i < mapnum; i++) diff --git a/libgomp/testsuite/libgomp.c/usm_env_handling-2.c b/libgomp/testsuite/libgomp.c/usm_env_handling-2.c new file mode 100644 index 00000000000..2683e1091c2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/usm_env_handling-2.c @@ -0,0 +1,157 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var GOMP_RUNTIME_USM "enabled" } */ + +/* Check that a global static variable is consistently updated on the device. + + This file checks that 'TARGET' data is handled properly when there is no + ALWAYS clause, i.e. that the device consistently accesses the device data + and not mixing host and device data. + + Actually, this is handled already by the compiler as such variables are + not replaced by a reference to the passed argument but directly access + the global variable. - But still useful to check to confirm that his works. */ + +#include +#include +#include +#include + +#ifdef USE_SELF_MAPS +#pragma omp requires self_maps + constexpr bool req_self_maps = true; +#else + constexpr bool req_self_maps = false; +#endif + +int a = 9; +int B[10] = {}; +struct S { + int x,y; +}; +struct S s = { .x = 1, .y = 2}; +struct S t = { .x = 3, .y = 4}; +#pragma omp declare target enter(a, B, s, t) + +int c[1] = {}; +int d = 0; +#pragma omp declare target enter(c, d) + + +void check_c_d_incr () +{ + if (c[0] != 42 + 3) + __builtin_abort (); + if (d != 27 + 4) + __builtin_abort (); + c[0] += 7; + d += 5; +} + +void check_vals_plus_inc3 () +{ + bool add_one = req_self_maps || omp_is_initial_device (); + if (a != 9 + 21 + (int)add_one) + __builtin_abort (); + for (int i = 0; i < 3; i++) + if (B[i] != 0 + 21 * i + (i == 3 && add_one)) + if (s.x != 1 + 21 + (int)add_one) + __builtin_abort (); + if (s.y != 2 + 21 + (int)add_one) + __builtin_abort (); + if (t.x != 3 + 21 + (int)add_one) + __builtin_abort (); + if (t.y != 4 + 21 + (int)add_one) + __builtin_abort (); + + a += 3; + for (int i = 0; i < 3; i++) + B[i] += 3 * i; + s.x += 3; + s.y += 3; + t.x += 3; + t.y += 3; +} + +int +main () +{ + // Debug output + const char *env_var = getenv ("GOMP_RUNTIME_USM"); + bool is_shared_mem = false; + #pragma omp target map(to: is_shared_mem) + is_shared_mem = true; + bool is_usm = (is_shared_mem && env_var + && strcasecmp ("enabled", env_var) == 0 + && omp_get_num_devices () > 0); + printf ("DEBUG: GOMP_RUNTIME_USM = %s, is_shared_mem = %s -> usm = %s\n", + env_var ? env_var : "", is_shared_mem ? "true" : "false", + is_usm ? "true" : "false"); + + + // Target update + modifying the data in TARGET + + c[0] = 42; + d = 27; + #pragma omp target update to(c, d) + + #pragma omp target map(present, alloc: d) + { + c[0] += 3; + d += 4; + } + + #pragma omp target + check_c_d_incr(); + + #pragma omp target update from(c, d) + if (c[0] != 42 + 3 + 7) + __builtin_abort (); + if (d != 27 + 4 + 5) + __builtin_abort (); + + // Some more tests involving structs and similar + + // Those use the initial value - unless the following + // affects the device value. + // That is only the case with 'requires self_maps' or + // when the device is the host (fallback) + + a++; + B[3]++; + s.x++; + s.y++; + t.x++; + t.y++; + + #pragma omp target map(from : a) /* to prevent firstprivate - no op as infinite ref count */ + { + a += 21; + for (int i = 0; i < 10; i++) + B[i] += 21*i; + s.x += 21; + s.y += 21; + t.x += 21; + t.y += 21; + } + + #pragma omp target + check_vals_plus_inc3 (); + + // Get the values from the device + #pragma omp target exit data map(always, from: a, s.x, B[0:5]) + #pragma omp target update from(s.y, B[5:5], t) + + bool add_one = req_self_maps || omp_get_num_devices () == 0; + if (a != 9 + 21 + 3 + (int)add_one) + __builtin_abort (); + for (int i = 0; i < 3; i++) + if (B[i] != 0 + (21 + 3) * i + (i == 3 && add_one)) + if (s.x != 1 + 21 + 3 + (int)add_one) + __builtin_abort (); + if (s.y != 2 + 21 + 3 + (int)add_one) + __builtin_abort (); + if (t.x != 3 + 21 + 3 + (int)add_one) + __builtin_abort (); + if (t.y != 4 + 21 + 3 + (int)add_one) + __builtin_abort (); +}