[Openmp-commits] [PATCH] D62393: [OPENMP][NVPTX]Mark parallel level counter as volatile.

[Johannes Doerfert via Phabricator via Openmp-commits]

jdoerfert added a subscriber: ye-luo.
jdoerfert added a comment.

I want to investigate the racy accesses further and make sure it is not a miscompile inside LLVM. 
I extracted the test case (see below) but I was not seeing the `ERROR`. How did you run the test case to see a different value for `Count`?

@ye-luo ran it with Clang-ykt (2019-05-27) on a Quadro P1000 <https://reviews.llvm.org/P1000> and I ran Clang (Jun 13 15:24:11 2019 : 3bc6e2a7aa3853b06045c42e81af094647c48676 <https://reviews.llvm.org/rG3bc6e2a7aa3853b06045c42e81af094647c48676>) on a V100 so far.

  #include <stdio.h>
  
  int main() {
  
    for (int i = 0; i < 1000; ++i) {
      int Count = 0;
  #pragma omp target parallel for reduction(+: Count) schedule(dynamic, 2) num_threads(64)
      for (int J = 0; J < 1000; ++J) {
        Count += J;
      }
      if (Count != 499500)
        printf("ERROR [@%i] %i\n", i, Count);
    }
  
    // Final result of Count is 1000 * (999-0) / 2
    // CHECK: Expected count with dynamic scheduling = 499500
    // printf("Expected count with dynamic scheduling = %d\n", Count);
  }


Repository:
  rOMP OpenMP

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

https://reviews.llvm.org/D62393