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