OpenMP offload Program with atomic does not finished when build with -O3 of AMD GPUs

[fel-cab]

On a MI250x

The following program doesn't finish when compiled as: clang -fopenmp -O3 --offload-arch=gfx90a

But it does finish when compiled as: clang -fopenmp --offload-arch=gfx90a

#include <stdio.h>
#include <stdlib.h>

#pragma omp requires atomic_default_mem_order(acq_rel)

int main() {

  int x = 0, y = 0;
  int errors = 0;

#pragma omp target parallel num_threads(2) map(tofrom: x, y, errors)
   {
       int thrd = omp_get_thread_num();
       if (thrd == 0) {
          x = 10;
          #pragma omp atomic write 
          y = 1;
       }
       if (thrd == 1){
          int tmp = 0;
          while (tmp == 0) {
            #pragma omp atomic read 
            tmp = y;
          }
          if(x != 10) errors++;
       }
   }
   if(errors)
     printf("Test Failed\n");
   else
     printf("Test Passed\n");

   return errors;
}

[llvmbot]

@llvm/issue-subscribers-openmp

Author: Felipe Cabarcas (fel-cab)

On a MI250x The following program doesn't finish when compiled as: `clang -fopenmp -O3 --offload-arch=gfx90a ` But it does finish when compiled as: `clang -fopenmp --offload-arch=gfx90a ` ```#include <omp.h> #include <stdio.h> #include <stdlib.h> #pragma omp requires atomic_default_mem_order(acq_rel) int main() { int x = 0, y = 0; int errors = 0; #pragma omp target parallel num_threads(2) map(tofrom: x, y, errors) { int thrd = omp_get_thread_num(); if (thrd == 0) { x = 10; #pragma omp atomic write y = 1; } if (thrd == 1){ int tmp = 0; while (tmp == 0) { #pragma omp atomic read tmp = y; } if(x != 10) errors++; } } if(errors) printf("Test Failed\n"); else printf("Test Passed\n"); return errors; } ```

[llvmbot]

@llvm/issue-subscribers-offload

Author: Felipe Cabarcas (fel-cab)

On a MI250x The following program doesn't finish when compiled as: `clang -fopenmp -O3 --offload-arch=gfx90a ` But it does finish when compiled as: `clang -fopenmp --offload-arch=gfx90a ` ```#include <omp.h> #include <stdio.h> #include <stdlib.h> #pragma omp requires atomic_default_mem_order(acq_rel) int main() { int x = 0, y = 0; int errors = 0; #pragma omp target parallel num_threads(2) map(tofrom: x, y, errors) { int thrd = omp_get_thread_num(); if (thrd == 0) { x = 10; #pragma omp atomic write y = 1; } if (thrd == 1){ int tmp = 0; while (tmp == 0) { #pragma omp atomic read tmp = y; } if(x != 10) errors++; } } if(errors) printf("Test Failed\n"); else printf("Test Passed\n"); return errors; } ```

[llvmbot]

@llvm/issue-subscribers-bug

Author: Felipe Cabarcas (fel-cab)

On a MI250x The following program doesn't finish when compiled as: `clang -fopenmp -O3 --offload-arch=gfx90a ` But it does finish when compiled as: `clang -fopenmp --offload-arch=gfx90a ` ```#include <omp.h> #include <stdio.h> #include <stdlib.h> #pragma omp requires atomic_default_mem_order(acq_rel) int main() { int x = 0, y = 0; int errors = 0; #pragma omp target parallel num_threads(2) map(tofrom: x, y, errors) { int thrd = omp_get_thread_num(); if (thrd == 0) { x = 10; #pragma omp atomic write y = 1; } if (thrd == 1){ int tmp = 0; while (tmp == 0) { #pragma omp atomic read tmp = y; } if(x != 10) errors++; } } if(errors) printf("Test Failed\n"); else printf("Test Passed\n"); return errors; } ```

[shiltian]

This is a backend issue.

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

Author: Felipe Cabarcas (fel-cab)

On a MI250x The following program doesn't finish when compiled as: `clang -fopenmp -O3 --offload-arch=gfx90a ` But it does finish when compiled as: `clang -fopenmp --offload-arch=gfx90a ` ```#include <omp.h> #include <stdio.h> #include <stdlib.h> #pragma omp requires atomic_default_mem_order(acq_rel) int main() { int x = 0, y = 0; int errors = 0; #pragma omp target parallel num_threads(2) map(tofrom: x, y, errors) { int thrd = omp_get_thread_num(); if (thrd == 0) { x = 10; #pragma omp atomic write y = 1; } if (thrd == 1){ int tmp = 0; while (tmp == 0) { #pragma omp atomic read tmp = y; } if(x != 10) errors++; } } if(errors) printf("Test Failed\n"); else printf("Test Passed\n"); return errors; } ```

[jdoerfert]

This code cannot reliably work. You are trying to synchronize within a warp on an AMDGPU, that does not work. OpenMP 6.0 makes this explicit. If you want to test atomics, split the reader and writer into different warps at least.