| // Use the generic state machine. On some architectures, other threads in the | |
| // main thread's warp must avoid barrier instructions. | |
| // | |
| // REQUIRES: gpu | |
| // RUN: %libomptarget-compile-run-and-check-generic | |
| // SPMDize. There is no main thread, so there's no issue. | |
| // | |
| // RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt > %t.spmd 2>&1 | |
| // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=SPMD -input-file=%t.spmd | |
| // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=SPMD -input-file=%t.spmd | |
| // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic | |
| // | |
| // SPMD: Transformed generic-mode kernel to SPMD-mode. | |
| // Use the custom state machine, which must avoid the same barrier problem as | |
| // the generic state machine. | |
| // | |
| // RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt \ | |
| // RUN: -Xoffload-linker -mllvm=-openmp-opt-disable-spmdization \ | |
| // RUN: -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&1 | |
| // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom | |
| // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom | |
| // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic | |
| // | |
| // Repeat with reduction clause, which has managed to break the custom state | |
| // machine in the past. | |
| // | |
| // RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt \ | |
| // RUN: -DADD_REDUCTION \ | |
| // RUN: -Xoffload-linker -mllvm=-openmp-opt-disable-spmdization \ | |
| // RUN: -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&1 | |
| // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom | |
| // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom | |
| // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic | |
| // https://github.com/llvm/llvm-project/issues/182119 | |
| // UNSUPPORTED: intelgpu | |
| // | |
| // CUSTOM: Rewriting generic-mode kernel with a customized state machine. | |
| #if ADD_REDUCTION | |
| #define REDUCTION(...) reduction(__VA_ARGS__) | |
| #else | |
| #define REDUCTION(...) | |
| #endif | |
| #include <stdio.h> | |
| int main() { | |
| int x = 0, y = 1; | |
| #pragma omp target teams num_teams(1) map(tofrom : x, y) REDUCTION(+ : x) | |
| { | |
| x += 5; | |
| #pragma omp parallel | |
| y = 6; | |
| } | |
| // CHECK: 5, 6 | |
| printf("%d, %d\n", x, y); | |
| return 0; | |
| } |