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.bc 2 // 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.out 3 4 // host-no-diagnostics 5 6 [[omp::assume("omp_no_openmp")]] void baz(void); 7 bar(void)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 foo(void)15void foo(void) { 16 #pragma omp target teams // #2 17 // 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 parallel 21 { 22 } 23 bar(); 24 #pragma omp parallel 25 { 26 } 27 } 28 } 29 spmd(void)30void spmd(void) { 31 // Verify we do not emit the remarks above for "SPMD" regions. 32 #pragma omp target teams 33 #pragma omp parallel 34 { 35 } 36 37 #pragma omp target teams distribute parallel for 38 for (int i = 0; i < 100; ++i) { 39 } 40 } 41 42 #pragma omp begin declare target device_type(nohost) 43 struct KernelEnvironmentTy; 44 struct KernelLaunchEnvironmentTy; 45 __attribute__((weak)) __kmpc_target_init(struct KernelEnvironmentTy *,struct KernelLaunchEnvironmentTy *)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 target 50 51 // expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}} 52