This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4.1] Support #pragma omp target {enter,exit} data
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: gcc-patches at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>
- Date: Thu, 30 Jul 2015 17:40:15 +0300
- Subject: Re: [gomp4.1] Support #pragma omp target {enter,exit} data
- Authentication-results: sourceware.org; auth=none
- References: <20150630125702 dot GI10247 at tucnak dot redhat dot com> <20150630154201 dot GB27446 at msticlxl57 dot ims dot intel dot com> <20150630161044 dot GM10247 at tucnak dot redhat dot com> <20150701210658 dot GA51887 at msticlxl57 dot ims dot intel dot com> <20150706153425 dot GA52133 at msticlxl57 dot ims dot intel dot com> <20150706172509 dot GY10247 at tucnak dot redhat dot com> <20150706184530 dot GB52133 at msticlxl57 dot ims dot intel dot com> <20150706204210 dot GB10247 at tucnak dot redhat dot com> <20150729190652 dot GA44830 at msticlxl57 dot ims dot intel dot com> <20150730081259 dot GM1780 at tucnak dot redhat dot com>
On Thu, Jul 30, 2015 at 10:12:59 +0200, Jakub Jelinek wrote:
> This test will fail on HSA, you don't assume just that it doesn't
> fallback to host, but also non-shared address space.
> I think it would be better to start with some check for non-shared address
> space, like:
> /* This test relies on non-shared address space. Punt otherwise. */
> void ensure_nonshared_as (void)
> {
> int a = 8;
> #pragma omp target map(to:a)
> {
> a++;
> }
> if (a == 8)
> exit (0);
> }
>
> And generally, it is better to have most of the tests not relying on
> offloading only or even non-shared address space, so that we also test
> shared address space and host fallback. But a few tests won't hurt...
Sure, but it's not possible to fully test data mapping without non-shared
address space. I've created new check_effective_target, ok for gomp-4_1-branch?
* testsuite/lib/libgomp.exp
(check_effective_target_offload_device_nonshared_as): New.
* testsuite/libgomp.c++/examples-4/e.53.2.C: Replace offload_device with
offload_device_nonshared_as.
* testsuite/libgomp.c/target-11.c: Ditto.
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 438777f..3a29b78 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -320,6 +320,19 @@ proc check_effective_target_offload_device { } {
} ]
}
+# Return 1 if offload device is available and it has non-shared address space.
+proc check_effective_target_offload_device_nonshared_as { } {
+ return [check_runtime_nocache offload_device_nonshared_as {
+ int main ()
+ {
+ int a = 8;
+ #pragma omp target map(to: a)
+ a++;
+ return a != 8;
+ }
+ } ]
+}
+
# Return 1 if at least one nvidia board is present.
proc check_effective_target_openacc_nvidia_accel_present { } {
diff --git a/libgomp/testsuite/libgomp.c++/examples-4/e.53.2.C b/libgomp/testsuite/libgomp.c++/examples-4/e.53.2.C
index 75276e7..6d5b5e4 100644
--- a/libgomp/testsuite/libgomp.c++/examples-4/e.53.2.C
+++ b/libgomp/testsuite/libgomp.c++/examples-4/e.53.2.C
@@ -1,5 +1,5 @@
// { dg-do run }
-// { dg-require-effective-target offload_device }
+// { dg-require-effective-target offload_device_nonshared_as }
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.c/target-11.c b/libgomp/testsuite/libgomp.c/target-11.c
index b86097a..ed6a17a 100644
--- a/libgomp/testsuite/libgomp.c/target-11.c
+++ b/libgomp/testsuite/libgomp.c/target-11.c
@@ -1,4 +1,4 @@
-/* { dg-require-effective-target offload_device } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
#include <stdlib.h>
#include <assert.h>
-- Ilya