[Openmp-commits] [PATCH] D104418: [PoC][WIP][OpenMP][Offloading] Fixed data race in libomptarget caused by async data movement

[Shilei Tian via Phabricator via Openmp-commits]

tianshilei1992 added a comment.

> Following your example, if T1 and T2 are really concurrent, I think it should be an application issue (data race), if they both perform initial mapping on the same memory without synchronization. I would bring this to the OpenMP LC.

Not exactly. Two concurrent target regions can read from same memory region.

  int array[1024];
  #pragma omp target map(to: data) nowait
  { /* read array */ }
  #pragma omp target map(to: data) nowait
  { /* read array */ }

I think this is a valid program, no matter from user's perspective or specification's perspective. How to deal with data mapping is implementation details IMO, but implementation should guarantee that `array` in the two regions should be same. It is not the case for us currently.

> For `bug49334.cpp` this means, that `depend(in: BlockA[0], BlockB[0])` is not sufficient and the dependency should rather be `depend(inout: BlockA[0], BlockB[0])`. The out dependency reflects the writing to the device copy of the variables or, respectively, the possible initial mapping.

Correct, but in the body of the task, `BlockA[0]` and `BlockB[0]` are only read. It's valid to mark them as `in` only.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D104418