* [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test
@ 2022-10-12 14:05 Tobias Burnus
2022-10-17 7:36 ` *ping* / " Tobias Burnus
2022-10-20 10:46 ` Jakub Jelinek
0 siblings, 2 replies; 3+ messages in thread
From: Tobias Burnus @ 2022-10-12 14:05 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Andrew Stubbs; +Cc: Vollweiler, Marcel
[-- Attachment #1: Type: text/plain, Size: 1201 bytes --]
This came up because the USM implementation with -foffload-memory={unified,pinned}
as posted at https://gcc.gnu.org/pipermail/gcc-patches/2022-July/597976.html
does not handle USM with static variables.
This shows up for the OG12 alias devel/omp/gcc-12 branch as FAIL for requires-4.c.
The attached patch prepares for skipping requires-4.c for the gcn/nvptx device
and adds an adjacent requires-4a.c testcase, using heap memory, that can still
run on gcn/nvptx.
Additionally, I commented on no longer used #defined, following the
precedence GOMP_DEVICE_HOST_NONSHM.
Thus, this tests adds another testcase and one effective-target check,
out-comments a unused #define - and that's it.
(Otherwise, it is just a prep patch.)
OK for mainline?
Tobias
PS: Currently, neither the preexisting offload_device_nvptx nor the new
offload_device_gcn target selector is used, neither in old code nor by this patch.
-----------------
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
[-- Attachment #2: libgomp-target-require.diff --]
[-- Type: text/x-patch, Size: 5428 bytes --]
libgomp: Add offload_device_gcn check, add requires-4a.c test
Duplicate libgomp.c-c++-common/requires-4.c (as ...-4a.c) but
with using a heap-allocated instead of static memory for a variable.
This change and the added offload_device_gcn check prepare for
pseudo-USM, where the device hardware cannot access all host
memory but only managed and pinned memory; for those, requires-4.c
will fail and the new check permits to add
target { ! { offload_device_nvptx || offload_device_gcn } }
to requires-4.c; however, it has not been added yet as pseuo-USM
support is not yet on mainline. (Review is pending for the USM
patches.)
include/ChangeLog:
* gomp-constants.h (GOMP_DEVICE_HSA): Comment (unused).
libgomp/ChangeLog:
* testsuite/lib/libgomp.exp (check_effective_target_offload_device_gcn):
New.
* testsuite/libgomp.c-c++-common/on_device_arch.h (device_arch_gcn,
on_device_arch_gcn): New.
* testsuite/libgomp.c-c++-common/requires-4a.c: New test; copied from
requires-4.c but using heap-allocated memory.
include/gomp-constants.h | 2 +-
libgomp/testsuite/lib/libgomp.exp | 12 +++++++
.../libgomp.c-c++-common/on_device_arch.h | 13 ++++++++
.../testsuite/libgomp.c-c++-common/requires-4a.c | 39 ++++++++++++++++++++++
4 files changed, 65 insertions(+), 1 deletion(-)
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 84316f953d0..fac7316b858 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -229,9 +229,9 @@ enum gomp_map_kind
/* #define GOMP_DEVICE_HOST_NONSHM 3 removed. */
#define GOMP_DEVICE_NOT_HOST 4
#define GOMP_DEVICE_NVIDIA_PTX 5
#define GOMP_DEVICE_INTEL_MIC 6
-#define GOMP_DEVICE_HSA 7
+/* #define GOMP_DEVICE_HSA 7 removed. */
#define GOMP_DEVICE_GCN 8
/* We have a compatibility issue. OpenMP 5.2 introduced
omp_initial_device with value of -1 which clashes with our
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 107a3c2ac9d..4b8c64de8a5 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -414,8 +414,20 @@ proc check_effective_target_offload_device_nvptx { } {
}
} ]
}
+# Return 1 if using a GCN offload device.
+proc check_effective_target_offload_device_gcn { } {
+ return [check_runtime_nocache offload_device_gcn {
+ #include <omp.h>
+ #include "testsuite/libgomp.c-c++-common/on_device_arch.h"
+ int main ()
+ {
+ return !on_device_arch_gcn ();
+ }
+ } ]
+}
+
# Return 1 if at least one Nvidia GPU is accessible.
proc check_effective_target_openacc_nvidia_accel_present { } {
return [check_runtime openacc_nvidia_accel_present {
diff --git a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
index f92743b04d7..6f66dbd784c 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
+++ b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
@@ -6,15 +6,22 @@ device_arch_nvptx (void)
{
return GOMP_DEVICE_NVIDIA_PTX;
}
+/* static */ int
+device_arch_gcn (void)
+{
+ return GOMP_DEVICE_GCN;
+}
+
/* static */ int
device_arch_intel_mic (void)
{
return GOMP_DEVICE_INTEL_MIC;
}
#pragma omp declare variant (device_arch_nvptx) match(construct={target},device={arch(nvptx)})
+#pragma omp declare variant (device_arch_gcn) match(construct={target},device={arch(gcn)})
#pragma omp declare variant (device_arch_intel_mic) match(construct={target},device={arch(intel_mic)})
/* static */ int
device_arch (void)
{
@@ -36,8 +43,14 @@ on_device_arch_nvptx ()
{
return on_device_arch (GOMP_DEVICE_NVIDIA_PTX);
}
+int
+on_device_arch_gcn ()
+{
+ return on_device_arch (GOMP_DEVICE_GCN);
+}
+
int
on_device_arch_intel_mic ()
{
return on_device_arch (GOMP_DEVICE_INTEL_MIC);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c
new file mode 100644
index 00000000000..4fb9783a97a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c
@@ -0,0 +1,39 @@
+/* { dg-additional-options "-flto" } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+/* { dg-additional-sources requires-4-aux.c } */
+
+/* Same as requires-4.c, but uses heap memory for 'a'. */
+
+/* Check no diagnostic by device-compiler's or host compiler's lto1.
+ Other file uses: 'requires reverse_offload', but that's inactive as
+ there are no declare target directives, device constructs nor device routines */
+
+/* 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.
+
+ No offload devices support USM at present, so we may verify host-fallback
+ execution by presence of separate memory spaces. */
+
+#pragma omp requires unified_address,unified_shared_memory
+
+int *a;
+extern void foo (void);
+
+int
+main (void)
+{
+ a = (int *) __builtin_calloc (sizeof (int), 10);
+ #pragma omp target map(to: a)
+ for (int i = 0; i < 10; i++)
+ a[i] = i;
+
+ for (int i = 0; i < 10; i++)
+ if (a[i] != i)
+ __builtin_abort ();
+
+ foo ();
+ __builtin_free (a);
+ return 0;
+}
^ permalink raw reply [flat|nested] 3+ messages in thread
* *ping* / Re: [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test
2022-10-12 14:05 [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test Tobias Burnus
@ 2022-10-17 7:36 ` Tobias Burnus
2022-10-20 10:46 ` Jakub Jelinek
1 sibling, 0 replies; 3+ messages in thread
From: Tobias Burnus @ 2022-10-17 7:36 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Andrew Stubbs; +Cc: Vollweiler, Marcel
On 12.10.22 16:05, Tobias Burnus wrote:
> This came up because the USM implementation with
> -foffload-memory={unified,pinned}
> as posted at
> https://gcc.gnu.org/pipermail/gcc-patches/2022-July/597976.html
> does not handle USM with static variables.
>
> This shows up for the OG12 alias devel/omp/gcc-12 branch as FAIL for
> requires-4.c.
>
> The attached patch prepares for skipping requires-4.c for the
> gcn/nvptx device
> and adds an adjacent requires-4a.c testcase, using heap memory, that
> can still
> run on gcn/nvptx.
>
> Additionally, I commented on no longer used #defined, following the
> precedence GOMP_DEVICE_HOST_NONSHM.
>
> Thus, this tests adds another testcase and one effective-target check,
> out-comments a unused #define - and that's it.
> (Otherwise, it is just a prep patch.)
>
> OK for mainline?
>
> Tobias
>
> PS: Currently, neither the preexisting offload_device_nvptx nor the new
> offload_device_gcn target selector is used, neither in old code nor by
> this patch.
-----------------
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
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test
2022-10-12 14:05 [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test Tobias Burnus
2022-10-17 7:36 ` *ping* / " Tobias Burnus
@ 2022-10-20 10:46 ` Jakub Jelinek
1 sibling, 0 replies; 3+ messages in thread
From: Jakub Jelinek @ 2022-10-20 10:46 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches, Andrew Stubbs, Vollweiler, Marcel
On Wed, Oct 12, 2022 at 04:05:32PM +0200, Tobias Burnus wrote:
> include/ChangeLog:
>
> * gomp-constants.h (GOMP_DEVICE_HSA): Comment (unused).
Comment out unused define.
or so, please.
> libgomp/ChangeLog:
>
> * testsuite/lib/libgomp.exp (check_effective_target_offload_device_gcn):
> New.
> * testsuite/libgomp.c-c++-common/on_device_arch.h (device_arch_gcn,
> on_device_arch_gcn): New.
> * testsuite/libgomp.c-c++-common/requires-4a.c: New test; copied from
> requires-4.c but using heap-allocated memory.
Otherwise LGTM.
Jakub
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2022-10-20 10:47 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-10-12 14:05 [Patch] libgomp: Add offload_device_gcn check, add requires-4a.c test Tobias Burnus
2022-10-17 7:36 ` *ping* / " Tobias Burnus
2022-10-20 10:46 ` Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).