public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
* [Bug libgomp/108624] New: [OpenMP] Map(ptr, ptr[2:N-2]) will cause a segfault, 'map(ptr, ptr[0:N])'
@ 2023-02-01  8:11 burnus at gcc dot gnu.org
  2023-02-01 11:54 ` [Bug libgomp/108624] " jules at gcc dot gnu.org
  0 siblings, 1 reply; 2+ messages in thread
From: burnus at gcc dot gnu.org @ 2023-02-01  8:11 UTC (permalink / raw)
  To: gcc-bugs

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=108624

            Bug ID: 108624
           Summary: [OpenMP] Map(ptr, ptr[2:N-2]) will cause a segfault,
                    'map(ptr, ptr[0:N])'
           Product: gcc
           Version: 13.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

Created attachment 54382
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=54382&action=edit
Testcase doing map(ptr1[1:N-1]) - segfaulting at runtime when doing offload

This came up during the OpenMP F2F and the segfaults at runtime both with GCC
and Clang,

It is based on
https://github.com/OpenMP/Examples/blob/main/devices/sources/target_ptr_map.1.c
- And we believe the modified variant should work (but I have not fully checked
whether it is indeed supposed to work).

The problem seems to be that
  ptr  maps the pointer
  ptr[n:m] maps the array

Without 'ptr', it is attached and firstprivatized - that just works fine.

With 'ptr' being present in addition, with
  ptr[0:m]  (i.e. starting at the lowest index)
the pointer attach of the data to the 'ptr' workes.

However, if one starts with
  ptr[2:N-2]  (i.e. with offset)
the pointer attachment does not work - leading to a segfault at runtime.

Naming remark:
  For map(ptr[2:N-2]),
    ptr[2] is in the mapped range
    while p[0] is in the extended range.

^ permalink raw reply	[flat|nested] 2+ messages in thread

* [Bug libgomp/108624] [OpenMP] Map(ptr, ptr[2:N-2]) will cause a segfault, 'map(ptr, ptr[0:N])'
  2023-02-01  8:11 [Bug libgomp/108624] New: [OpenMP] Map(ptr, ptr[2:N-2]) will cause a segfault, 'map(ptr, ptr[0:N])' burnus at gcc dot gnu.org
@ 2023-02-01 11:54 ` jules at gcc dot gnu.org
  0 siblings, 0 replies; 2+ messages in thread
From: jules at gcc dot gnu.org @ 2023-02-01 11:54 UTC (permalink / raw)
  To: gcc-bugs

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=108624

jules at gcc dot gnu.org changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |jules at gcc dot gnu.org

--- Comment #1 from jules at gcc dot gnu.org ---
This works with the in-review series:

https://gcc.gnu.org/pipermail/gcc-patches/2022-December/609031.html

i.e. for the mapping:

#pragma omp target map(ptr1[1:N-1]) map(ptr2[:N])

we get (gimplify):

_1 = ptr1 + 4;
map(tofrom:*_1 [len: 252]) map(firstprivate:ptr1 [pointer assign, bias: 4])

and for the mapping:

#pragma omp target map(ptr1, ptr1[1:N-1]) map(ptr2[:N])

we get (gimplify):

_4 = ptr1.2_3 + 4
map(tofrom:ptr1 [len: 8]) map(tofrom:*_4 [len: 252]) map(attach:ptr1 [bias: 4])

the test runs fine either way.

^ permalink raw reply	[flat|nested] 2+ messages in thread

end of thread, other threads:[~2023-02-01 11:54 UTC | newest]

Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-02-01  8:11 [Bug libgomp/108624] New: [OpenMP] Map(ptr, ptr[2:N-2]) will cause a segfault, 'map(ptr, ptr[0:N])' burnus at gcc dot gnu.org
2023-02-01 11:54 ` [Bug libgomp/108624] " jules at gcc dot gnu.org

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).