start of array is mapped regardless of explicit map clause

[jdenny-ornl]

I'm seeing unexpected behavior starting at @ye-luo's recent c1a6fe1 (D123093): the start of an array is mapped regardless of its explicit map clause. I can still reproduce at today's 82e5976. I've tested while offloading to nvptx64.

For example, the following prints 1, but it used to print 0 before c1a6fe1:

#include <omp.h>
#include <stdio.h>
int main(int argc, char *argv[]) {
  int arr[1000];
  #pragma omp target data map(tofrom:arr[998:1])
  {
    printf("%d\n", omp_target_is_present(arr, omp_get_default_device()));
  }
  return 0;
}

[llvmbot]

@llvm/issue-subscribers-openmp

[ye-luo]

Since spec 5.1

If the value of a given pointer is in the mapped address range of a currently mapped list item then 
that currently mapped list item is a matching mapped list item. Otherwise, if the value of the
pointer is in the extended address range of a currently mapped list item then that currently mapped
list item is a matching mapped list item.

In your case, the pointer lands in the extended address range. So it is considered found in the existing maps.
I'd like to see more opinions. Maybe this can be refined.

[jdenny-ornl]

Looking at OpenMP 5.2 sec. 18.8.3 "omp_target_is_present", "Effect":

The omp_target_is_present routine returns true if device_num refers to the host device or
if ptr refers to storage that has corresponding storage in the device data environment of device
device_num. Otherwise, the routine returns false.

In my example, arr does not refer to storage that has corresponding storage in the device data environment, right?

[jdenny-ornl]

Keep in mind that my example uses a target data construct. The first sentence of the section you're quoting says it talks about target constructs:

This section describes how a pointer that is predetermined firstprivate for a target construct may
be assigned an initial value that is the address of an object that exists in a device data environment
and corresponds to a matching mapped list item.

That doesn't apply in my example, right?

[ye-luo]

I think you are right. I need to find a way to distinguish both cases.

[RaviNarayanaswamy]

Even though target data maps 1 element of arr at offset 998., the present check on arr should return true.

[jdenny-ornl]

Even though target data maps 1 element of arr at offset 998., the present check on arr should return true.

Are you arguing that arr does "refer to storage that has corresponding storage in the device data environment"? Or does the omp_target_is_present description need to change?

[RaviNarayanaswamy]

What I am saying is that arr has a corresponding storage on the device. Not all memory of host is present on device and is an error accessing them. Can only access memory which is allocated on device

[ye-luo]

omp_target_is_present, returns true if

1. pointer found in the mapped address range
2. pointer found in the extended address range

If I understood correctly.
Jeol is saying 1), Ravi is saying 2)
I see pros and cons of either choices. Let me ask the language committee.

[jdenny-ornl]

Surely map(present,alloc:arr[0:1]) is still expected to produce a runtime error after map(alloc:arr[998:1]). I suppose I expected omp_target_is_present to use the same definition of "present".

Yes, please let us know what the language committee says.

Thanks for everyone's replies.

[ye-luo]

Surely map(present,alloc:arr[0:1]) is still expected to produce a runtime error after map(alloc:arr[998:1]). I suppose I expected omp_target_is_present to use the same definition of "present".

Yes, please let us know what the language committee says.

Thanks for everyone's replies.

Indeed, I did exactly the same test on the present clause. There is no problem because there is no ambiguity.
arr[begin:size] array notation carries 3 pieces of info. Base address, begin address and size(end address).
omp_target_is_present(arr) only sees the base address. It is not equivalent to map(present,alloc:arr[0:1]) which checks 4 bytes (int) instead of 1.

[jdenny-ornl]

I understand omp_target_is_present doesn't take a size, so I assumed it would check the presence of a single byte. That brings more questions to mind:

• What should omp_get_mapped_ptr(arr, dev) return? I'm assuming NULL, but maybe it's like omp_target_is_present instead?
• What should omp_target_is_accessible(arr, 1, dev) return? Surely false.
• What should omp_target_is_accessible(arr, 0, dev) return?

I realize omp_get_mapped_ptr and omp_target_is_accessible are not implemented yet upstream. I had prototyped them in my work and was planning to contribute that, but now I'm suspicious I've implemented them incorrectly.

[RaviNarayanaswamy]

What should omp_get_mapped_ptr(arr, dev) return?
return the associated arr pointer on the device
What should omp_target_is_accessible(arr, 1, dev) return? Surely false.
Yes false
What should omp_target_is_accessible(arr, 0, dev) return?
true

[ye-luo]

• omp_get_mapped_ptr(arr, dev) the pointer lookup should succeed. So return the associated arr pointer on the device.
• omp_target_is_accessible tests whether host memory is accessible from a given device. It has nothing to do with mapping. It is unclear to me what is the intention and how to implement this. In CUDA/HIP, even if you have pinned host memory. the point may not be used for accessing host on the device and you need cudaHostGetDevicePointer to translate to a pointer accessible on the device. However, I don't know if the translated pointer is still accessible from host. The host pointer and the translated pointer point to the same memory but the address may or may not be different.

[jdenny-ornl]

What should omp_target_is_accessible(arr, 0, dev) return?
true

Where does the spec clarify the size=0 case?

omp_target_is_accessible tests whether host memory is accessible from a given device. It has nothing to do with mapping.

Ravi's answer above for omp_target_is_accessible(arr, 0, dev) suggests it does see mappings. I assume omp_target_is_accessible(&arr[998], 1, dev) would also return true.

Somewhere I heard that omp_target_is_accessible returns true for unified shared memory as well as mapped memory, and omp_target_is_present returns true only for mapped memory. Maybe that's wrong. I don't trust my own read of the spec here.

[RaviNarayanaswamy]

The way I look at it is that if you specify a size for omp_target_is_accessible then the runtime checks if that size is allocated on the device then return true, if the size is zero, then you then return true if there is corresponding pointer on the device.

[jdenny-ornl]

The way I look at it is that if you specify a size for omp_target_is_accessible then the runtime checks if that size is allocated on the device then return true, if the size is zero, then you then return true if there is corresponding pointer on the device.

That sounds reasonable to me, but I cannot figure out where the spec says these things.

[ye-luo]

Added a patch https://reviews.llvm.org/D123891