[flang-commits] [flang] [llvm] [Flang][OpenMP] Fix Fortran pointer map(present) null handling (PR #204615)

[Abhinav Gaba via flang-commits]

================
@@ -0,0 +1,34 @@
+! REQUIRES: flang, amdgpu
+! RUN: %libomptarget-compile-fortran-generic
+! RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
+
+program map_present_pointer
+  implicit none
+  integer, target :: src(4) = [10, 20, 30, 40]
+  integer, pointer :: p(:) => null()
+  integer :: out
+
+  out = -1
+!$omp target map(present, to: p) map(tofrom: out)
----------------
abhinavgaba wrote:

Expecting this present check to not result in runtime error is not OpenMP compliant.

This has previously been discussed in the spec committee. The `map(present:p[0])` check is similar to the an `assert(omp_target_is_present(p, ...)`, and if `p` is `null`, then it's expected to fail.

Here's an equivalent example in C:

```c
❯ cat tgt_present_nullptr.c
#include <omp.h>
#include <stdio.h>

int main() {
  int *p = (void *)0;
  printf("omp_target_is_present: %d\n",
         omp_target_is_present(p, omp_get_default_device()));
#pragma omp target map(present, alloc : p[0 : 0])
  {
    printf("%p\n", p);
  }
}
❯ clang -O0 -fopenmp -fopenmp-targets=x86_64 tgt_present_nullptr.c && ./a.out
omp_target_is_present: 0
omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x0000000000000000 (0 bytes)
omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
omptarget error: Call to targetDataBegin failed, abort target.
omptarget error: Failed to process data before launching the kernel.
omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
omptarget error: Source location information not present. Compile with -g or -gline-tables-only.
omptarget fatal error 1: failure of target construct while offloading is mandatory
[1]    224786 IOT instruction (core dumped)  ./a.out
``` 

This is different from how OpenACC's `present` works. There is an open issue (https://github.com/OpenMP/spec/issues/3544#issuecomment-1796284774) to allow something like `accessible_or_null` to OpenMP to align with OpenACC's `present`, but it's not part of 6.1. It might be something worth considering for 7.0.

https://github.com/llvm/llvm-project/pull/204615