[Openmp-commits] [PATCH] D134268: [Clang][OpenMP] Codegen generation for has_device_addr claues.

[Shilei Tian via Phabricator via Openmp-commits]

tianshilei1992 added a comment.

Hi there, I'm trying to fix https://github.com/llvm/llvm-project/issues/59160. The faulty case is basically like following:

  cpp
  void xoo() {
    short a[10], b[10];
    a[1] = 111;
    b[1] = 111;
  #pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b)
  #pragma omp target has_device_addr(a) has_device_addr(b[0])
    {
      a[1] = 222;
      b[1] = 222;
      // CHECK: 222 222
      printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b);
    }
    // CHECK:111
    printf("%hd %hd %p %p %p\n", a[1], b[1], &a, b, &b); // 111 111 p1d p2d p3d
  }

It looks like at runtime, we are trying to copy a (device) pointer to a device pointer by using host to device data transfer. I noticed that's because we have `TO` flag marked for the argument. However, since `a` and `b` are in `has_device_addr`, we are not supposed to map the two variables right?



================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:9098
+      CombinedInfo.Types.push_back(
+          (Cap->capturesVariable() ? OMP_MAP_TO : OMP_MAP_LITERAL) |
+          OMP_MAP_TARGET_PARAM);
----------------
The variable used in `has_device_addr` indicates it already has device address and can be accessed directly. Why do we need `OMP_MAP_TO` here? `OMP_MAP_TO` indicates we are gonna transfer data to the device right?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D134268