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

Shilei Tian via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Thu Jan 12 11:49:48 PST 2023

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:

  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) |
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?

  rG LLVM Github Monorepo



More information about the Openmp-commits mailing list