[Openmp-commits] [PATCH] D68100: [OpenMP 5.0] declare mapper runtime implementation

[George Rokos via Phabricator via Openmp-commits]

grokos added a comment.

@lildmh: I think I've found a bug. I used `declare_mapper_target.cpp`. When we call the mapper function, it generates 3 components. The first 2 are identical and correspond to the parent struct. This is what MapperComponents looks like inside function `target_data_mapper`:

  (gdb) print MapperComponents
  $1 = {Components = std::vector of length 3, capacity 4 = {{Base = 0x7fffffffb598, Begin = 0x7fffffffb598, Size = 8, Type = 32}, {Base = 0x7fffffffb598, Begin = 0x7fffffffb598, Size = 8, Type = 32}, {Base = 0x7fffffffb598,
        Begin = 0x62efd0, Size = 4096, Type = 562949953421331}}}

Mapping the parent struct twice is problematic. If we have more struct members and some of them are NOT pointers, then upon `target_data_end` libomptarget will check the parent struct's reference counter to determine whether the scalar member must be copied back to the host. If the reference counter is greater than 1, then the runtime will skip copying back the scalar. Mapping the parent struct two times in a row results in `RefCount=2`.

So in the example below (modified `declare_mapper_target.cpp`) the scalar is processed by libomptarget but because at that time the struct's RefCount is 2 we never copy the scalar back:

  #include <cstdio>
  #include <cstdlib>
  #include <omp.h>
  
  #define NUM 1024
  
  class C {
  public:
    int *a;
    int onHost;
  };
  
  #pragma omp declare mapper(id: C s) map(s.a[0:NUM], s.onHost)
  
  int main() {
    C c;
    c.a = (int*) malloc(sizeof(int)*NUM);
    c.onHost = -1;
    for (int i = 0; i < NUM; i++) {
      c.a[i] = 1;
    }
    #pragma omp target teams distribute parallel for map(mapper(id),tofrom: c)
    for (int i = 0; i < NUM; i++) {
      ++c.a[i];
      if (i == 0) {
        c.onHost = omp_is_initial_device();
      }
    }
  
    int sum = 0;
    for (int i = 0; i < NUM; i++) {
      sum += c.a[i];
    }
    printf("Executed on %s\n", c.onHost==1 ? "host" : c.onHost==0 ? "device" : "unknown");
    // CHECK: Sum = 2048
    printf("Sum = %d\n", sum);
    return 0;
  }

Upon `target_data_end` the mapper function will generate this:

  (gdb) print MapperComponents
  $1 = {Components = std::vector of length 4, capacity 4 = {{Base = 0x7fffffffb588, Begin = 0x7fffffffb588, Size = 16, Type = 32}, {Base = 0x7fffffffb588, Begin = 0x7fffffffb588, Size = 12, Type = 32}, {Base = 0x7fffffffb588,
        Begin = 0x62efd0, Size = 4096, Type = 562949953421331}, {Base = 0x7fffffffb588, Begin = 0x7fffffffb590, Size = 4, Type = 562949953421315}}}

When libomptarget processes the scalar, the parent struct's RefCount is 2, so inside the if-block in `omptarget.cpp:507-524` `CopyMember` will never be set to true and the scalar will never be copied back to the host.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D68100