Hi Tobias! On 2022-07-07T11:36:34+0200, Tobias Burnus wrote: > On 07.07.22 10:42, Thomas Schwinge wrote: >> In preparation for other changes: > ... >> On 2022-06-29T16:33:02+0200, Tobias Burnus wrote: >>> +/* { dg-output "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" } */ >> (The latter diagnostic later got conditionalized by 'GOMP_DEBUG=1'.) >> OK to push the attached "Enhance 'libgomp.c-c++-common/requires-4.c', >> 'libgomp.c-c++-common/requires-5.c' testing"? > ... >> libgomp/ >> * testsuite/libgomp.c-c++-common/requires-4.c: Enhance testing. >> * testsuite/libgomp.c-c++-common/requires-5.c: Likewise. > ... >> --- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c >> +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c >> @@ -1,22 +1,29 @@ >> -/* { dg-do link { target offloading_enabled } } */ >> /* { dg-additional-options "-flto" } */ >> /* { dg-additional-sources requires-4-aux.c } */ >> >> -/* Check diagnostic by device-compiler's or host compiler's lto1. >> +/* Check no diagnostic by device-compiler's or host compiler's lto1. > > I note that without ENABLE_OFFLOADING that there is never any lto1 > diagnostic. > > However, given that no diagnostic is expected, it also works for "! > offloading_enabled". > > Thus, the change fine. ACK. >> Other file uses: 'requires reverse_offload', but that's inactive as >> there are no declare target directives, device constructs nor device routines */ >> >> +/* For actual offload execution, prints the following (only) if GOMP_DEBUG=1: >> + "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" >> + and does host-fallback execution. */ > > The latter is only true when also device code is produced – and a device > is available for that/those device types. I think that's what you imply > by "For actual offload execution" ACK. > but it is a bit hidden. > > Maybe s/For actual offload execution, prints/It may print/ is clearer? I've settled on: /* Depending on offload device capabilities, it may print something like the following (only) if GOMP_DEBUG=1: "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" and in that case does host-fallback execution. */ > In principle, it would be nice if we could test for the output, but > currently setting an env var for remote execution does not work, yet. > Cf. https://gcc.gnu.org/pipermail/gcc-patches/2022-July/597773.html Right, I'm aware of that issue with remote testing, and that's why I didn't propose such output verification. (In a few other test cases, we do have 'dg-set-target-env-var GOMP_DEBUG "1"', which then at present are UNSUPPORTED for remote testing.) > When set, we could use offload_target_nvptx etc. (..._amdgcn, ..._any) > to test – as this guarantees that it is compiled for that device + the > device is available. Use 'target offload_device_nvptx', not 'target offload_target_nvptx', etc. ;-) >> + >> #pragma omp requires unified_address,unified_shared_memory >> >> -int a[10]; >> +int a[10] = { 0 }; >> extern void foo (void); >> >> int >> main (void) >> { >> - #pragma omp target >> + #pragma omp target map(to: a) > > Hmm, I wonder whether I like it or not. Without, there is an implicit > "map(tofrom:a)". On the other hand, OpenMP permits that – even with > unified-shared memory – the implementation my copy the data to the > device. (For instance, to permit faster access to "a".) > > Thus, ... > >> + for (int i = 0; i < 10; i++) >> + a[i] = i; >> + >> for (int i = 0; i < 10; i++) >> - a[i] = 0; >> + if (a[i] != i) >> + __builtin_abort (); > ... this condition (back on the host) could also fail with USM. However, > given that to my knowledge no USM implementation actually copies the > data, I believe it is fine. Right, this is meant to describe/test the current GCC master branch behavior, where USM isn't supported, so I didn't consider that. But I agree, a source code comment should be added: As no offload devices support USM at present, we may verify host-fallback execution by absence of separate memory spaces. */ > (Disclaimer: I have not checked what OG12, > but I guess it also does not copy it.) >> --- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c >> +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c >> @@ -1,21 +1,25 @@ >> -/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn } } } */ >> /* { dg-additional-sources requires-5-aux.c } */ >> >> +/* For actual offload execution, prints the following (only) if GOMP_DEBUG=1: >> + "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" >> + and does host-fallback execution. */ >> + > This wording is correct with the now-removed check – but if you remove > the offload_target..., it only "might" print it, depending, well, on the > conditions set by offload_target... >> #pragma omp requires unified_shared_memory, unified_address, reverse_offload >> >> -int a[10]; >> +int a[10] = { 0 }; >> extern void foo (void); >> >> int >> main (void) >> { >> - #pragma omp target >> + #pragma omp target map(to: a) >> + for (int i = 0; i < 10; i++) >> + a[i] = i; >> + >> for (int i = 0; i < 10; i++) >> - a[i] = 0; >> + if (a[i] != i) >> + __builtin_abort (); >> >> foo (); >> return 0; >> } > > Thus: LGTM, if you update the GOMP_DEBUG=... wording, either using > "might" (etc.) or by being more explicit. Used the wording as in 'libgomp.c-c++-common/requires-4.c'. Thanks for the review! Pushed to master branch commit 5647e2c3853cbd51a6536a84b8eb0eb3c210dfbf "Enhance 'libgomp.c-c++-common/requires-4.c', 'libgomp.c-c++-common/requires-5.c' testing", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955