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 <complex> 11 12 // Verify we do not add tbaa metadata to type punned memory operations: 13 14 // CHECK: call i64 @__kmpc_shuffle_int64( 15 // CHECK-NEXT: store i64 %{{.*}}, i64* %{{.*}}, align {{[0-9]+$}} 16 17 // CHECK: call i64 @__kmpc_shuffle_int64( 18 // CHECK-NEXT: store i64 %{{.*}}, i64* %{{.*}}, align {{[0-9]+$}} 19 20 template <typename T> 21 void complex_reduction() { 22 #pragma omp target teams distribute 23 for (int ib = 0; ib < 100; ib++) { 24 std::complex<T> partial_sum; 25 const int istart = ib * 4; 26 const int iend = (ib + 1) * 4; 27 #pragma omp parallel for reduction(+ \ 28 : partial_sum) 29 for (int i = istart; i < iend; i++) 30 partial_sum += std::complex<T>(i, i); 31 } 32 } 33 34 void test() { 35 complex_reduction<float>(); 36 complex_reduction<double>(); 37 } 38 #endif 39