[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#1679193 <https://reviews.llvm.org/D62393#1679193>, @protze.joachim wrote:

> It is probably not directly related to this patch, but I think the logic of `omp_in_parallel()` is not valid. It only works, because the level counting is also broken. This code
>
>   int main(){
>   #pragma omp target 
>     #pragma omp parallel if(0)
>       if(omp_get_thread_num()==0) {
>         printf("omp_in_parallel=%i\n", omp_in_parallel()); 
>         printf("omp_get_level=%i\n", omp_get_level()); 
>         printf("omp_get_active_level=%i\n", omp_get_active_level()); 
>       }
>     return 0;
>   }
>
>
> should print (and does so with `env OMP_TARGET_OFFLOAD=disabled`):
>
>   omp_in_parallel=0
>   omp_get_level=1
>   omp_get_active_level=0
>
>
> As I understand the code for `omp_in_parallel()`, the example code would print `omp_in_parallel=1` if the level would correctly be increased.


You're wrong. My output:

  omp_in_parallel=0
  omp_get_level=1
  omp_get_active_level=0

So, `omp_in_parallel` works correctly on GPU.

If it does not work correctly with `OMP_TARGET_OFFLOAD=disabled`, the problem is definitely in `libomp`, not `libomptarget`. With disabled offloading the code is not offloaded to GPU.


Repository:
  rOMP OpenMP

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

https://reviews.llvm.org/D62393