[Bug libgomp/82194] New: Mapping array section (e.g. [0:N-1]) using omp target map crashes at runtime

["josem at udel dot edu" <gcc-bugzilla at gcc dot gnu dot org>]

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

            Bug ID: 82194
           Summary: Mapping array section (e.g. [0:N-1]) using omp target
                    map crashes at runtime
           Product: gcc
           Version: unknown
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: josem at udel dot edu
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

Created attachment 42161
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=42161&action=edit
Reproducible for array_segment bug

Hello GCC community, 

--- SUMARY:

This bug affects OpenMP 4.5 implementation. It occurs when trying to map an
array section by using the [lower-bound:length] subscript expressions.

--- MORE INFORMATION:

According to the OpenMP 4.5 standard, section 2.4, page 44, it is possible to
map an array section containing a subset of the original elements of the array.
However, I am getting a runtime error when trying to do so. Attached you will
find an example code that crashes when trying to use map(from: a1d[1:]) which
should map the array from element 1 to N.

Output error is: 
libgomp: Trying to map into device [0x3fffffc66a6c..0x3fffffc67a0c) object when
[0x3fffffc66a70..0x3fffffc67a0c) is already mapped

Other options that fail as well with the same error (different memory regions)
are: 
> #pragma omp target data map(from: a1d[1:N-1]) 
> #pragma omp target data map(from: a1d[0:N-1])
> #pragma omp target data map(from: a1d[0:N/5])

Backtrace information: 
#0  0x00003fffb7d2847c in write () from /lib64/libc.so.6
#1  0x00003fffb7ca7274 in _IO_new_file_write () from /lib64/libc.so.6
#2  0x00003fffb7ca7da8 in __GI__IO_file_xsputn () from /lib64/libc.so.6
#3  0x00003fffb7c76c50 in buffered_vfprintf () from /lib64/libc.so.6
#4  0x00003fffb7c77170 in vfprintf@@GLIBC_2.17 () from /lib64/libc.so.6
#5  0x00003fffb7e4bb2c in gomp_verror (fmt=0x3fffb7e78d98 "Trying to map into
device [%p..%p) object when [%p..%p) is already mapped", list=0x3fffffffc198
"\374\305\377\377\377?")
    at ../../../gcc/libgomp/error.c:62
#6  0x00003fffb7e4bbe8 in gomp_vfatal (fmt=<optimized out>, list=0x3fffffffc198
"\374\305\377\377\377?") at ../../../gcc/libgomp/error.c:79
#7  0x00003fffb7e4bc38 in gomp_fatal (fmt=<optimized out>) at
../../../gcc/libgomp/error.c:89
#8  0x00003fffb7e63dd0 in gomp_map_vars_existing (devicep=0x3fffb7e698f4
<GOMP_target_ext+1716>, oldn=0x10514a70, newn=<optimized out>, newn=<optimized
out>, kind=<optimized out>, tgt_var=0x10514c00)
    at ../../../gcc/libgomp/target.c:234
#9  0x00003fffb7e67450 in gomp_map_vars_existing (newn=0x3fffffffc288,
newn=0x3fffffffc288, kind=<optimized out>, tgt_var=0x10514c00, oldn=0x10514a70,
devicep=<optimized out>)
    at ../../../gcc/libgomp/target.c:234
#10 gomp_map_vars (devicep=0x100cb500, mapnum=2, hostaddrs=0x3fffffffc5e8,
devaddrs=0x0, sizes=0x10080020 <.omp_data_sizes.8.3210>, kinds=0x10080030
<.omp_data_kinds.9.3211>, short_mapkind=true,
    pragma_kind=GOMP_MAP_VARS_TARGET) at ../../../gcc/libgomp/target.c:725
#11 0x00003fffb7e698f4 in GOMP_target_ext (device=<optimized out>,
fn=0x10000bcc <fail_on_map._omp_fn.0>, mapnum=2, hostaddrs=0x3fffffffc5e8,
sizes=0x10080020 <.omp_data_sizes.8.3210>,
    kinds=0x10080030 <.omp_data_kinds.9.3211>, flags=<optimized out>,
depend=<optimized out>, args=0x3fffffffc5c8) at
../../../gcc/libgomp/target.c:2029
#12 0x00000000100009b8 in fail_on_map ()
#13 0x0000000010000b34 in main ()

--- COMPILATION COMMAND: 

> gcc -std=c99 -fopenmp -foffload="-lm" -lm array_segment_map.c -o array_segment_map.c.o

--- MACHINE INFORMATION

I tested this in a 64 bit system. 
> gcc --version
gcc (GCC) 7.1.1 20170718
- OS: Red Hat Enterprise Linux Server 7.3 - CentOS release 7
- PROCESSOR: POWER8NVL ppc64le
- TARGET DEVICE: NVIDIA Tesla P100 GPU