diff mbox series

Add 'libgomp.c/address-space-1.c' (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref')

Message ID 871r6kqhot.fsf@dem-tschwing-1.ger.mentorg.com
State New
Headers show
Series Add 'libgomp.c/address-space-1.c' (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') | expand

Commit Message

Thomas Schwinge Aug. 23, 2021, 3:55 p.m. UTC 
Hi!

On 2021-08-20T16:49:25+0200, Jakub Jelinek <jakub@redhat.com> wrote:
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.c/address-space-1.c
>> @@ -0,0 +1,24 @@
>> +/* Verify OMP instances of variables with address space.  */
>> +
>> +/* { dg-do run { target i?86-*-* x86_64-*-* } } */
>> +/* { dg-require-effective-target offload_device_nonshared_as } */

I also added:

    +/* With Intel MIC (emulated) offloading:
    +       offload error: process on the device 0 unexpectedly exited with code 0
    +   { dg-xfail-run-if TODO { offload_device_intel_mic } } */

Might this be a symptom related to my earlier comment:

| [...] -- shouldn't we force generic
| address space for all 'tree' types read in via LTO streaming for
| offloading compilation?  I assume that (in the general case) address
| spaces are never compatible between host and offloading compilation?
| For the attached "Add 'libgomp.c/address-space-1.c'", propagating the
| '__seg_fs' address space across the offloading boundary (assuming I did
| interpret the dumps correctly) doesn't seem to cause any problems, but
| maybe it's problematic for other cases?

..., or it's yet another problem, of course.  ;-)

>> +#include <assert.h>
>> +
>> +int __seg_fs a;
>> +
>> +int
>> +main (void)
>> +{
>> +  // a = 123; // SIGSEGV
>> +  int b;
>> +#pragma omp target map(alloc: a) map(from: b)
>> +  {
>> +    a = 321; // no SIGSEGV (given 'offload_device_nonshared_as')
>> +    asm volatile ("" : : : "memory");
>
> Maybe better asm volatile ("" : : "g" (&a) : "memory");
> so that the compiler doesn't think it could optimize it away to
> just b = 321;

Thanks.

>> +    b = a;
>> +  }
>> +  assert (b == 321);
>> +
>> +  return 0;
>> +}

Pushed "Add 'libgomp.c/address-space-1.c'" to master branch in
commit 29c355f76ceeb4639c21acaf52c50d35c8472720, 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
diff mbox series

Patch

From 29c355f76ceeb4639c21acaf52c50d35c8472720 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu, 19 Aug 2021 15:14:51 +0200
Subject: [PATCH] Add 'libgomp.c/address-space-1.c'

Intel MIC (emulated) offloading execution failure remains to be analyzed.

	libgomp/
	* testsuite/libgomp.c/address-space-1.c: New file.

Co-authored-by: Jakub Jelinek <jakub@redhat.com>
---
 libgomp/testsuite/libgomp.c/address-space-1.c | 28 +++++++++++++++++++
 1 file changed, 28 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.c/address-space-1.c

diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c
new file mode 100644
index 00000000000..6ad57deec42
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/address-space-1.c
@@ -0,0 +1,28 @@ 
+/* Verify OMP instances of variables with address space.  */
+
+/* { dg-do run { target i?86-*-* x86_64-*-* } } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* With Intel MIC (emulated) offloading:
+       offload error: process on the device 0 unexpectedly exited with code 0
+   { dg-xfail-run-if TODO { offload_device_intel_mic } } */
+
+#include <assert.h>
+
+int __seg_fs a;
+
+int
+main (void)
+{
+  // a = 123; // SIGSEGV
+  int b;
+#pragma omp target map(alloc: a) map(from: b)
+  {
+    a = 321; // no SIGSEGV (given 'offload_device_nonshared_as')
+    asm volatile ("" : : "g" (&a) : "memory");
+    b = a;
+  }
+  assert (b == 321);
+
+  return 0;
+}
-- 
2.25.1