52 lines · c
1// RUN: %clang_cc1 -verify=host -Rpass=openmp-opt -Rpass-analysis=openmp-opt -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc2// RUN: %clang_cc1 -verify -Rpass=openmp-opt -Rpass-analysis=openmp-opt -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out3 4// host-no-diagnostics5 6[[omp::assume("omp_no_openmp")]] void baz(void);7 8void bar(void) {9#pragma omp parallel // #1 \10 // expected-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}11 {12 }13}14 15void foo(void) {16#pragma omp target teams // #217 // expected-remark@#2 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}18 {19 baz(); // expected-remark {{Value has potential side effects preventing SPMD-mode execution. Add `[[omp::assume("ompx_spmd_amenable")]]` to the called function to override. [OMP121]}}20#pragma omp parallel21 {22 }23 bar();24#pragma omp parallel25 {26 }27 }28}29 30void spmd(void) {31 // Verify we do not emit the remarks above for "SPMD" regions.32#pragma omp target teams33#pragma omp parallel34 {35 }36 37#pragma omp target teams distribute parallel for38 for (int i = 0; i < 100; ++i) {39 }40}41 42#pragma omp begin declare target device_type(nohost)43struct KernelEnvironmentTy;44struct KernelLaunchEnvironmentTy;45__attribute__((weak))46extern "C" int __kmpc_target_init(struct KernelEnvironmentTy *, struct KernelLaunchEnvironmentTy *) { // expected-remark {{Could not internalize function. Some optimizations may not be possible. [OMP140]}}47 return 0;48}49#pragma omp end declare target50 51// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}}52