offload/test/offloading/bug51781.c

   1 // Use the generic state machine.  On some architectures, other threads in the
   2 // main thread's warp must avoid barrier instructions.
   3 //
   4 // RUN: %libomptarget-compile-run-and-check-generic
   5
   6 // SPMDize.  There is no main thread, so there's no issue.
   7 //
   8 // RUN: %libomptarget-compile-generic -O2 -Rpass=openmp-opt > %t.spmd 2>&1
   9 // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=SPMD -input-file=%t.spmd
  10 // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=SPMD -input-file=%t.spmd
  11 // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
  12 //
  13 // SPMD: Transformed generic-mode kernel to SPMD-mode.
  14
  15 // Use the custom state machine, which must avoid the same barrier problem as
  16 // the generic state machine.
  17 //
  18 // RUN: %libomptarget-compile-generic -O2 -Rpass=openmp-opt \
  19 // RUN:   -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&1
  20 // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom
  21 // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom
  22 // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
  23 //
  24 // Repeat with reduction clause, which has managed to break the custom state
  25 // machine in the past.
  26 //
  27 // RUN: %libomptarget-compile-generic -O2 -Rpass=openmp-opt -DADD_REDUCTION \
  28 // RUN:   -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&1
  29 // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom
  30 // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom
  31 // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
  32 //
  33 // CUSTOM: Rewriting generic-mode kernel with a customized state machine.
  34
  35 #if ADD_REDUCTION
  36 #define REDUCTION(...) reduction(__VA_ARGS__)
  37 #else
  38 #define REDUCTION(...)
  39 #endif
  40
  41 #include <stdio.h>
  42 int main() {
  43   int x = 0, y = 1;
  44 #pragma omp target teams num_teams(1) map(tofrom : x, y) REDUCTION(+ : x)
  45   {
  46     x += 5;
  47 #pragma omp parallel
  48     y = 6;
  49   }
  50   // CHECK: 5, 6
  51   printf("%d, %d\n", x, y);
  52   return 0;
  53 }