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 -O1 -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 -O1 -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 -O1 -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 // XFAIL: amdgcn-amd-amdhsa
  35
  36 #if ADD_REDUCTION
  37 #define REDUCTION(...) reduction(__VA_ARGS__)
  38 #else
  39 #define REDUCTION(...)
  40 #endif
  41
  42 #include <stdio.h>
  43 int main() {
  44   int x = 0, y = 1;
  45 #pragma omp target teams num_teams(1) map(tofrom : x, y) REDUCTION(+ : x)
  46   {
  47     x += 5;
  48 #pragma omp parallel
  49     y = 6;
  50   }
  51   // CHECK: 5, 6
  52   printf("%d, %d\n", x, y);
  53   return 0;
  54 }