openmp/libomptarget/deviceRTLs/nvptx/test/parallel/flush.c

   1 // RUN: %compile-run-and-check
   2
   3 #include <omp.h>
   4 #include <stdio.h>
   5
   6 int main(int argc, char *argv[]) {
   7   int data, out, flag = 0;
   8 #pragma omp target parallel num_threads(64) map(tofrom                         \
   9                                                 : out, flag) map(to            \
  10                                                                  : data)
  11   {
  12     if (omp_get_thread_num() == 0) {
  13       /* Write to the data buffer that will be read by thread */
  14       data = 42;
  15 /* Flush data to thread 32 */
  16 #pragma omp flush(data)
  17       /* Set flag to release thread 32 */
  18 #pragma omp atomic write
  19       flag = 1;
  20     } else if (omp_get_thread_num() == 32) {
  21       /* Loop until we see the update to the flag */
  22       int val;
  23       do {
  24 #pragma omp atomic read
  25         val = flag;
  26       } while (val < 1);
  27       out = data;
  28 #pragma omp flush(out)
  29     }
  30   }
  31   // CHECK: out=42.
  32   /* Value of out will be 42 */
  33   printf("out=%d.\n", out);
  34   return !(out == 42);
  35 }