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-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-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-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 internal void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}_worker() [[attr0:#[0-9]*]] 15 // CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}() [[attr0]] 16 // CHECK: %call = call float @_Z3sinf(float 0.000000e+00) [[attr5:#[0-9]*]] 17 // CHECK-DAG: declare i32 @llvm.nvvm.read.ptx.sreg.warpsize() [[attr1:#[0-9]*]] 18 // CHECK-DAG: declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() [[attr1]] 19 // CHECK-DAG: declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() [[attr1]] 20 // CHECK: declare void @__kmpc_kernel_init(i32, i16) 21 // CHECK-NOT: # 22 // CHECK: declare void @__kmpc_data_sharing_init_stack() 23 // CHECK-NOT: # 24 // CHECK: declare float @_Z3sinf(float) [[attr2:#[0-9]*]] 25 // CHECK: declare void @__kmpc_kernel_deinit(i16) 26 // CHECK-NOT: # 27 // CHECK: declare void @__kmpc_barrier_simple_spmd(%struct.ident_t*, i32) [[attr3:#[0-9]*]] 28 // CHECK: declare i1 @__kmpc_kernel_parallel(i8**) 29 // CHECK-NOT: # 30 // CHECK: declare i32 @__kmpc_global_thread_num(%struct.ident_t*) [[attr4:#[0-9]*]] 31 // CHECK: declare void @__kmpc_kernel_end_parallel() 32 // CHECK-NOT: # 33 // CHECK: define internal void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}_worker() [[attr0]] 34 // CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}() [[attr0]] 35 // CHECK: %call = call double @_Z3sind(double 0.000000e+00) [[attr5]] 36 // CHECK: declare double @_Z3sind(double) [[attr2]] 37 38 // CHECK: attributes [[attr0]] 39 // CHECK-NOT: "llvm.assume" 40 // CHECK: attributes [[attr1]] 41 // CHECK-NOT: "llvm.assume" 42 // CHECK: attributes [[attr2]] 43 // CHECK-SAME: "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations" 44 // CHECK: attributes [[attr3]] 45 // CHECK-NOT: "llvm.assume" 46 // CHECK: attributes [[attr4]] 47 // CHECK-NOT: "llvm.assume" 48 // CHECK: attributes [[attr5]] 49 // CHECK-SAME: "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations" 50 51 52 template <typename T> 53 void foo() { 54 cos(T(0)); 55 } 56 57 template <typename T> 58 void complex_reduction() { 59 foo<T>(); 60 #pragma omp target 61 sin(T(0)); 62 } 63 64 #pragma omp assumes ext_check_that_this_is_attached_to_included_functions_and_template_instantiations 65 66 void test() { 67 complex_reduction<float>(); 68 complex_reduction<double>(); 69 } 70 #endif 71