https://gcc.gnu.org/g:b590dc3fdc0aeb9bcd923cfd22778a0864102101
commit r17-3858-gb590dc3fdc0aeb9bcd923cfd22778a0864102101 Author: Tobias Burnus <[email protected]> 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. Diff: --- 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 f688f33378d6..081701ca6342 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 000000000000..2683e1091c21 --- /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 <omp.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> + +#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 : "<unset>", 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 (); +}
