1 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 2 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s 3 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 4 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s 5 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -fexceptions -fcxx-exceptions -aux-triple powerpc64le-unknown-unknown -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s 6 // expected-no-diagnostics 7 #ifndef HEADER 8 #define HEADER 9 10 #include <cmath> 11 12 // TODO: Think about teaching the OMPIRBuilder about default attributes as well so the __kmpc* declarations are annotated. 13 14 // CHECK: define weak_odr protected ptx_kernel void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}({{.*}}) [[attr0:#[0-9]]] 15 // CHECK: call i32 @__kmpc_target_init( 16 // CHECK: declare noundef float @_Z3sinf(float noundef) [[attr1:#[0-9]*]] 17 // CHECK: declare void @__kmpc_target_deinit( 18 // CHECK: define weak_odr protected ptx_kernel void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}({{.*}}) [[attr0]] 19 // CHECK: %call = call noundef double @_Z3sind(double noundef 0.000000e+00) [[attr2:#[0-9]]] 20 // CHECK: declare noundef double @_Z3sind(double noundef) [[attr1]] 21 22 // CHECK: attributes [[attr0]] 23 // CHECK-NOT: "llvm.assume" 24 // CHECK: attributes [[attr1]] 25 // CHECK-SAME: "llvm.assume"="ompx_check_that_this_is_attached_to_included_functions_and_template_instantiations" 26 // CHECK: attributes [[attr2]] 27 // CHECK-SAME: "llvm.assume"="ompx_check_that_this_is_attached_to_included_functions_and_template_instantiations" 28 29 30 template <typename T> 31 void foo() { 32 cos(T(0)); 33 } 34 35 template <typename T> 36 void complex_reduction() { 37 foo<T>(); 38 #pragma omp target 39 sin(T(0)); 40 } 41 42 #pragma omp assumes ext_check_that_this_is_attached_to_included_functions_and_template_instantiations 43 44 void test() { 45 complex_reduction<float>(); 46 complex_reduction<double>(); 47 } 48 #endif 49