Hi!
On 2021-08-20T16:49:25+0200, Jakub Jelinek <ja...@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
>From 29c355f76ceeb4639c21acaf52c50d35c8472720 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@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 <ja...@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
