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

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev added a comment.

In D62393#1542471 <https://reviews.llvm.org/D62393#1542471>, @jdoerfert wrote:

> I want to investigate the racy accesses further and make sure it is not a miscompile inside LLVM.


This is not a problem inside LLVM. The problem  appears after optimizations performed by the ptxas tool (when it compiles PTX to SASS) at O3 <https://reviews.llvm.org/owners/package/3/> with the inlined runtime.

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

You need to compile it with the inlined runtime at O2 <https://reviews.llvm.org/owners/package/2/> or O3 <https://reviews.llvm.org/owners/package/3/>. And you need the latest version of the libomptarget, clang-ykt version is implemented differently and has no this problem.

> @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