[OPENMP][NVPTX]Perform memory flush if number of threads to sync is 1 or less.

[ABataev]

Summary

According to the OpenMP standard, barrier operation must perform
implicit flush operation. Currently, if there is only one thread in the
team, barrier does not flush the memory. Patch fixes this problem.

[ABataev]

Addressed comments.

[ABataev]

Updated test.

[ABataev]

Rebase

[grokos]

This looks consistent with the changes in D62397, I think it's good to go.

libomptarget/deviceRTLs/nvptx/src/sync.cu

if (numberOfActiveOMPThreads > 1) {
             WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE);
 
         PRINT(LD_SYNC,
               "call kmpc_barrier with %d omp threads, sync parameter %d\n",
               (int)numberOfActiveOMPThreads, (int)threads);
         // Barrier #1 is for synchronization among active threads.
         named_sync(L1_BARRIER, threads);
       }
+    } else {
+      // Still need to flush the memory per the standard.
+      __kmpc_flush(loc_ref);
     } // numberOfActiveOMPThreads > 1
     PRINT0(LD_SYNC, "completed kmpc_barrier\n");
   }
 }
 
 // Emit a simple barrier call in SPMD mode.  Assumes the caller is in an L0
 // parallel region and that all worker threads participate.
 EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) {

libomptarget/deviceRTLs/nvptx/test/parallel/barrier.c

+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+  int data, flag = 0;
+#pragma omp target teams num_teams(2) map(tofrom : data, flag) thread_limit(1)
+#pragma omp parallel num_threads(1)
+  {
+    if (omp_get_team_num() == 0) {
+      /* Write to the data buffer that will be read by thread */
+      data = 42;
+/* Flush data to thread 1 and strictly order the write to data
+   relative to the write to the flag */
+#pragma omp barrier
+      /* Set flag to release thread 1 */
+      flag = 1;
+/* Flush flag to ensure that thread 1 sees the change */
+#pragma omp barrier
+    } else if (omp_get_team_num() == 1) {
+/* Loop until we see the update to the flag */
+#pragma omp barrier
+      while (flag < 1) {
+#pragma omp barrier
+      }
+#pragma omp barrier
+    }
+  }
+  // CHECK: data=42.
+  /* Value of data will be 42, value of flag still undefined */
+  printf("data=%d.\n", data);
+  return 0;
+}