https://gcc.gnu.org/bugzilla/show_bug.cgi?id=92854
--- Comment #4 from Thomas Schwinge <tschwinge at gcc dot gnu.org> --- Created attachment 47444 --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=47444&action=edit [WIP] 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c' (In reply to jules from comment #2) > For strictly-paired acc_map_data/acc_unmap_data calls that don't interfere > with other mappings -- no, probably not. But (like I guess you noticed too) > the existing code feels wrong (or at least ugly) nonetheless. ;-) > It's probably possible to come up with a legal (but odd) test case Odd is not the problem -- actual users are doing such things, too. > in which > the condition is false -- maybe something like this (untested!): > > int arr1[100], arr2[100], arr3[100]; > > void foo (void) > { > #pragma acc data copy(arr1) copy(arr2) copy(arr3) > { > void *darr2 = acc_deviceptr (arr2); > acc_map_data (arr2, darr2, sizeof (arr2)); > [...] > acc_unmap_data (arr2); > } > } > > Now I think the acc_map_data call will reuse the tgt for the structured data > mapping, which binds together several array copies into a single device > block. Forcing the first key's refcount to infinity (as is done in the > current implementation of acc_map_data) is also wrong in that case. In these cases, 'acc_map_data' will (should) refuse because "address [...] is already mapped", see the test cases I'm attaching. (ACK?) Can you easily (please don't spend much time on this) tell what's wrong with 'libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c' (see TODOs therein). If not, then we shall file a separate PR for this OpenACC 'declare' case -- GCC's OpenACC 'declare' implementation has many other issues, after all.
