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

Johannes Doerfert via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Thu Jun 13 13:00:15 PDT 2019

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

  rOMP OpenMP



More information about the Openmp-commits mailing list