public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] libgomp: Add no-target-region rev offload test + fix plugin-nvptx
@ 2022-11-28 14:29 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-11-28 14:29 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:97c07e0b841ac919c5eb6f4b8c335ede1f32de78
commit 97c07e0b841ac919c5eb6f4b8c335ede1f32de78
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Nov 28 15:16:47 2022 +0100
libgomp: Add no-target-region rev offload test + fix plugin-nvptx
OpenMP permits that a 'target device(ancestor:1)' is called without being
enclosed in a target region - using the current device (i.e. the host) in
that case. This commit adds a testcase for this.
In case of nvptx, the missing on-device 'GOMP_target_ext' call causes that
it and also the associated on-device GOMP_REV_OFFLOAD_VAR variable are not
linked in from nvptx's libgomp.a. Thus, handle the failing cuModuleGetGlobal
gracefully by disabling reverse offload and assuming that the failure is fine.
libgomp/ChangeLog:
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_load_image): Use unsigned int
for 'i' to match 'fn_entries'; regard absent GOMP_REV_OFFLOAD_VAR
as valid and the code having no reverse-offload code.
* testsuite/libgomp.c-c++-common/reverse-offload-2.c: New test.
(cherry picked from commit 9f9d128f459e0c5ace8f7b85504d277b5a838daf)
Diff:
---
libgomp/ChangeLog.omp | 10 +++++
libgomp/plugin/plugin-nvptx.c | 36 ++++++++++------
.../libgomp.c-c++-common/reverse-offload-2.c | 49 ++++++++++++++++++++++
3 files changed, 83 insertions(+), 12 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index b259d72e576..3360c9d7608 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,13 @@
+2022-11-28 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2022-11-25 Tobias Burnus <tobias@codesourcery.com>
+
+ * plugin/plugin-nvptx.c (GOMP_OFFLOAD_load_image): Use unsigned int
+ for 'i' to match 'fn_entries'; regard absent GOMP_REV_OFFLOAD_VAR
+ as valid and the code having no reverse-offload code.
+ * testsuite/libgomp.c-c++-common/reverse-offload-2.c: New test.
+
2022-11-28 Tobias Burnus <tobias@codesourcery.com>
Backported from master:
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index be3af6cea2a..d4c6dd119fe 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1423,7 +1423,8 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
else if (rev_fn_table)
{
CUdeviceptr var;
- size_t bytes, i;
+ size_t bytes;
+ unsigned int i;
r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &var, &bytes, module,
"$offload_func_table");
if (r != CUDA_SUCCESS)
@@ -1446,12 +1447,11 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
if (rev_fn_table && *rev_fn_table && dev->rev_data == NULL)
{
- /* cuMemHostAlloc memory is accessible on the device, if unified-shared
- address is supported; this is assumed - see comment in
- nvptx_open_device for CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING. */
- CUDA_CALL_ASSERT (cuMemHostAlloc, (void **) &dev->rev_data,
- sizeof (*dev->rev_data), CU_MEMHOSTALLOC_DEVICEMAP);
- CUdeviceptr dp = (CUdeviceptr) dev->rev_data;
+ /* Get the on-device GOMP_REV_OFFLOAD_VAR variable. It should be
+ available but it might be not. One reason could be: if the user code
+ has 'omp target device(ancestor:1)' in pure hostcode, GOMP_target_ext
+ is not called on the device and, hence, it and GOMP_REV_OFFLOAD_VAR
+ are not linked in. */
CUdeviceptr device_rev_offload_var;
size_t device_rev_offload_size;
CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal,
@@ -1459,11 +1459,23 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
&device_rev_offload_size, module,
XSTRING (GOMP_REV_OFFLOAD_VAR));
if (r != CUDA_SUCCESS)
- GOMP_PLUGIN_fatal ("cuModuleGetGlobal error - GOMP_REV_OFFLOAD_VAR: %s", cuda_error (r));
- r = CUDA_CALL_NOCHECK (cuMemcpyHtoD, device_rev_offload_var, &dp,
- sizeof (dp));
- if (r != CUDA_SUCCESS)
- GOMP_PLUGIN_fatal ("cuMemcpyHtoD error: %s", cuda_error (r));
+ {
+ free (*rev_fn_table);
+ *rev_fn_table = NULL;
+ }
+ else
+ {
+ /* cuMemHostAlloc memory is accessible on the device, if
+ unified-shared address is supported; this is assumed - see comment
+ in nvptx_open_device for CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING. */
+ CUDA_CALL_ASSERT (cuMemHostAlloc, (void **) &dev->rev_data,
+ sizeof (*dev->rev_data), CU_MEMHOSTALLOC_DEVICEMAP);
+ CUdeviceptr dp = (CUdeviceptr) dev->rev_data;
+ r = CUDA_CALL_NOCHECK (cuMemcpyHtoD, device_rev_offload_var, &dp,
+ sizeof (dp));
+ if (r != CUDA_SUCCESS)
+ GOMP_PLUGIN_fatal ("cuMemcpyHtoD error: %s", cuda_error (r));
+ }
}
nvptx_set_clocktick (module, dev);
diff --git a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-2.c b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-2.c
new file mode 100644
index 00000000000..33bd38481bb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-2.c
@@ -0,0 +1,49 @@
+/* { dg-do run } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+
+#pragma omp requires reverse_offload
+
+int
+main ()
+{
+ int A[10];
+ int y;
+
+ for (int i = 0; i < 10; i++)
+ A[i] = 2*i;
+
+ y = 42;
+
+ /* Pointlessly copy to the default device. */
+ #pragma omp target data map(to: A)
+ {
+ /* Not enclosed in a target region (= i.e. running on the host); the
+ following is valid - it runs on the current device (= host). */
+ #pragma omp target device ( ancestor:1 ) firstprivate(y) map(to: A)
+ {
+ if (y != 42)
+ __builtin_abort ();
+ for (int i = 0; i < 10; i++)
+ if (A[i] != 2*i)
+ __builtin_abort ();
+ for (int i = 0; i < 10; i++)
+ if (A[i] != 2*i)
+ A[i] = 4*i;
+ y = 31;
+ }
+
+ if (y != 42)
+ __builtin_abort ();
+ for (int i = 0; i < 10; i++)
+ if (A[i] != 2*i)
+ __builtin_abort ();
+ }
+
+ if (y != 42)
+ __builtin_abort ();
+ for (int i = 0; i < 10; i++)
+ if (A[i] != 2*i)
+ __builtin_abort ();
+
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-11-28 14:29 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-11-28 14:29 [gcc/devel/omp/gcc-12] libgomp: Add no-target-region rev offload test + fix plugin-nvptx Tobias Burnus
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).