| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex "__omp_rtl_" |
| // Test target codegen - host bc file has to be created first. |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-target-debug=111 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-EQ |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-DEFAULT |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-threads-oversubscription -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-THREADS |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-teams-oversubscription -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-TEAMS |
| // expected-no-diagnostics |
| |
| #ifndef HEADER |
| #define HEADER |
| |
| //. |
| // CHECK: @__omp_rtl_debug_kind = weak_odr constant i32 1 |
| // CHECK: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0 |
| // CHECK: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0 |
| //. |
| // CHECK-EQ: @__omp_rtl_debug_kind = weak_odr constant i32 111 |
| // CHECK-EQ: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0 |
| // CHECK-EQ: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0 |
| //. |
| // CHECK-DEFAULT: @__omp_rtl_debug_kind = weak_odr constant i32 0 |
| // CHECK-DEFAULT: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0 |
| // CHECK-DEFAULT: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0 |
| //. |
| // CHECK-THREADS: @__omp_rtl_debug_kind = weak_odr constant i32 0 |
| // CHECK-THREADS: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 0 |
| // CHECK-THREADS: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 1 |
| //. |
| // CHECK-TEAMS: @__omp_rtl_debug_kind = weak_odr constant i32 0 |
| // CHECK-TEAMS: @__omp_rtl_assume_teams_oversubscription = weak_odr constant i32 1 |
| // CHECK-TEAMS: @__omp_rtl_assume_threads_oversubscription = weak_odr constant i32 0 |
| //. |
| void foo() { |
| #pragma omp target |
| { } |
| } |
| |
| #endif |