57 lines · c
1// Use the generic state machine. On some architectures, other threads in the2// main thread's warp must avoid barrier instructions.3//4// RUN: %libomptarget-compile-run-and-check-generic5 6// SPMDize. There is no main thread, so there's no issue.7//8// RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt > %t.spmd 2>&19// RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=SPMD -input-file=%t.spmd10// RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=SPMD -input-file=%t.spmd11// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic12//13// SPMD: Transformed generic-mode kernel to SPMD-mode.14 15// Use the custom state machine, which must avoid the same barrier problem as16// the generic state machine.17//18// RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt \19// RUN: -Xoffload-linker -mllvm=-openmp-opt-disable-spmdization \20// RUN: -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&121// RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom22// RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom23// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic24//25// Repeat with reduction clause, which has managed to break the custom state26// machine in the past.27//28// RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt \29// RUN: -DADD_REDUCTION \30// RUN: -Xoffload-linker -mllvm=-openmp-opt-disable-spmdization \31// RUN: -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&132// RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom33// RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom34// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic35//36// CUSTOM: Rewriting generic-mode kernel with a customized state machine.37 38#if ADD_REDUCTION39#define REDUCTION(...) reduction(__VA_ARGS__)40#else41#define REDUCTION(...)42#endif43 44#include <stdio.h>45int main() {46 int x = 0, y = 1;47#pragma omp target teams num_teams(1) map(tofrom : x, y) REDUCTION(+ : x)48 {49 x += 5;50#pragma omp parallel51 y = 6;52 }53 // CHECK: 5, 654 printf("%d, %d\n", x, y);55 return 0;56}57