49 lines · cpp
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.bc2// 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 %s3// 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.bc4// 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 %s5// 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 %s6// expected-no-diagnostics7#ifndef HEADER8#define HEADER9 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 30template <typename T>31void foo() {32 cos(T(0));33}34 35template <typename T>36void complex_reduction() {37 foo<T>();38#pragma omp target39 sin(T(0));40}41 42#pragma omp assumes ext_check_that_this_is_attached_to_included_functions_and_template_instantiations43 44void test() {45 complex_reduction<float>();46 complex_reduction<double>();47}48#endif49