[PATCH] D141627: [Clang][OpenMP] Fix the issue that list items in `has_device_addr` are still mapped to the target device

[Abhinav Gaba via Phabricator via cfe-commits]

abhinavgaba added subscribers: dreachem, kkwli0.
abhinavgaba added a comment.

> In target data we already put a and b in use_device_addr. That indicates all use of a and b will be the corresponding device addresses. Therefore, in target directive, we should use is_device_address instead of has_device_addr. The correct way to use has_device_addr is, we already map the list items by using target data w/o use_device_addr. Then when we launch a kernel using target directive with has_device_addr, we tell the target region, the list items *should* be there, and use them, otherwise it is an error (we choose to error out for the undefined behavior).

I think you are talking about `is_device_ptr` clause. There is no `is_device_address` clause in OpenMP.

The is_device_ptr clause is meant only for "ptrs" (pointers). For example:

  int *p = omp_target_alloc(...);
  #pragma omp target is_device_ptr(p)



> On the other hand, has_device_addr indicates that the list items *should* have device address, which means there has to be an entry for that.

Based on a brief discussion with some members of the OpenMP spec committee, the idea for "has_device_addr" is to have the address passed-in directly (as a literal, similar to `is_device_ptr`) into the target region, without any map lookup. So, there is no requirement that the variable has to be mapped, or tracked by libomptarget. That requirement is for `map(present:x)`.

One example use-case from @dreachem is this:

  #pragma omp requires unified_shared_memory
  
  int x ;
  
  printg("%p\n", &x); // p1h
  printf("%d\n", omp_target_is_present(&x, omp_get_default_device())); // 0 (x is not "present", as per the OpenMP runtime)
  
  #pragma omp target has_device_addr(x)
  print("%p\n", &x); // p1h (same address as on the host side)
  }

In this case, because of unified shared memory, `x` is accessible on device as well, even though it is not mapped, or made declare target etc. So we need to pass the address of x into the region, even though `omp_target_is_present` would return false for it.

In terms of the code emitted, the original idea of passing the address in as a LITERAL, similar to `is_device_ptr` is the right way to think about it.

  %x = alloca i32 ; Original allocation for x
  
  Map:
  <ptr %x, ptr %x, sizeof(ptr), PARAM|LITERAL> 
  
  Outlined function:
  define void @outlined...(ptr %x) {
     ...
     call i32 @printf(..., ptr %x)
     ...
  }


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D141627/new/

https://reviews.llvm.org/D141627