1 // Test host codegen. 2 // RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128 3 // RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 4 // RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128 5 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST 6 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 7 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST 8 9 // Test target codegen - host bc file has to be created first. 10 // RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 11 // RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes=CHECK,INT128 12 // RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s 13 // RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,INT128 14 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 15 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s 16 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s 17 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s 18 19 // expected-no-diagnostics 20 #ifndef HEADER 21 #define HEADER 22 23 // HOST: @[[MAPTYPES_PRIVATE:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] 24 // HOST: @[[MAPTYPES_FIRSTPRIVATE:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] 25 // HOST: @[[MAPTYPES_REDUCTION:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35] 26 // HOST: @[[MAPTYPES_FROM:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 34] 27 // HOST: @[[MAPTYPES_TO:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 33] 28 // HOST: @[[MAPTYPES_ALLOC:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 32] 29 // HOST: @[[MAPTYPES_ARRAY_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 35, i64 35, i64 35] 30 // HOST: @[[MAPTYPES_ARRAY_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 33, i64 33, i64 33] 31 // HOST-INT128: @[[MAPTYPES_INT128_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 35, i64 35, i64 35] 32 // HOST-INT128: @[[MAPTYPES_INT128_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 34, i64 34, i64 34] 33 // 34 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_PRIVATE:__omp_offloading_[^"\\]*mapWithPrivate[^"\\]*]]\00" 35 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FIRSTPRIVATE:__omp_offloading_[^"\\]*mapWithFirstprivate[^"\\]*]]\00" 36 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_REDUCTION:__omp_offloading_[^"\\]*mapWithReduction[^"\\]*]]\00" 37 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FROM:__omp_offloading_[^"\\]*mapFrom[^"\\]*]]\00" 38 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_TO:__omp_offloading_[^"\\]*mapTo[^"\\]*]]\00" 39 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ALLOC:__omp_offloading_[^"\\]*mapAlloc[^"\\]*]]\00" 40 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R0:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00" 41 // CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R1:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00" 42 // INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R0:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00" 43 // INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R1:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00" 44 45 // HOST: define {{.*}}mapWithPrivate 46 // HOST: call {{.*}} @.[[OFFLOAD_PRIVATE]].region_id{{.*}} @[[MAPTYPES_PRIVATE]] 47 // 48 // CHECK: define {{.*}} void @[[OFFLOAD_PRIVATE]]() 49 // CHECK: call void ({{.*}}@[[OUTLINE_PRIVATE:.omp_outlined.[.0-9]*]] 50 // 51 // CHECK: define {{.*}} void @[[OUTLINE_PRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid.) 52 void mapWithPrivate() { 53 int x, y; 54 #pragma omp target teams private(x) map(x,y) private(y) 55 ; 56 } 57 58 // HOST: define {{.*}}mapWithFirstprivate 59 // HOST: call {{.*}} @.[[OFFLOAD_FIRSTPRIVATE]].region_id{{.*}} @[[MAPTYPES_FIRSTPRIVATE]] 60 // 61 // CHECK: define {{.*}} void @[[OFFLOAD_FIRSTPRIVATE]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y) 62 // CHECK: call void ({{.*}}@[[OUTLINE_FIRSTPRIVATE:.omp_outlined.[.0-9]*]] 63 // 64 // CHECK: define {{.*}} void @[[OUTLINE_FIRSTPRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x, i{{[0-9]*}} %y) 65 void mapWithFirstprivate() { 66 int x, y; 67 #pragma omp target teams firstprivate(x) map(x,y) firstprivate(y) 68 ; 69 } 70 71 // HOST: define {{.*}}mapWithReduction 72 // HOST: call {{.*}} @.[[OFFLOAD_REDUCTION]].region_id{{.*}} @[[MAPTYPES_REDUCTION]] 73 // 74 // CHECK: define {{.*}} void @[[OFFLOAD_REDUCTION]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y) 75 // CHECK: call void ({{.*}}@[[OUTLINE_REDUCTION:.omp_outlined.[.0-9]*]] 76 // 77 // CHECK: define {{.*}} void @[[OUTLINE_REDUCTION]]({{[^)]*}} i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y) 78 // CHECK: %.omp.reduction.red_list = alloca [2 x i8*] 79 void mapWithReduction() { 80 int x, y; 81 #pragma omp target teams reduction(+:x) map(x,y) reduction(+:y) 82 ; 83 } 84 85 // HOST: define {{.*}}mapFrom 86 // HOST: call {{.*}} @.[[OFFLOAD_FROM]].region_id{{.*}} @[[MAPTYPES_FROM]] 87 // 88 // CHECK: define {{.*}} void @[[OFFLOAD_FROM]](i{{[0-9]*}}* {{[^,]*}}%x) 89 // CHECK: call void ({{.*}}@[[OUTLINE_FROM:.omp_outlined.[.0-9]*]] 90 // 91 // CHECK: define {{.*}} void @[[OUTLINE_FROM]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x) 92 void mapFrom() { 93 int x; 94 #pragma omp target teams firstprivate(x) map(from:x) 95 ; 96 } 97 98 // HOST: define {{.*}}mapTo 99 // HOST: call {{.*}} @.[[OFFLOAD_TO]].region_id{{.*}} @[[MAPTYPES_TO]] 100 // 101 // CHECK: define {{.*}} void @[[OFFLOAD_TO]](i{{[0-9]*}}* {{[^,]*}}%x) 102 // CHECK: call void ({{.*}}@[[OUTLINE_TO:.omp_outlined.[.0-9]*]] 103 // 104 // CHECK: define {{.*}} void @[[OUTLINE_TO]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x) 105 void mapTo() { 106 int x; 107 #pragma omp target teams firstprivate(x) map(to:x) 108 ; 109 } 110 111 // HOST: define {{.*}}mapAlloc 112 // HOST: call {{.*}} @.[[OFFLOAD_ALLOC]].region_id{{.*}} @[[MAPTYPES_ALLOC]] 113 // 114 // CHECK: define {{.*}} void @[[OFFLOAD_ALLOC]](i{{[0-9]*}}* {{[^,]*}}%x) 115 // CHECK: call void ({{.*}}@[[OUTLINE_ALLOC:.omp_outlined.[.0-9]*]] 116 // 117 // CHECK: define {{.*}} void @[[OUTLINE_ALLOC]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x) 118 void mapAlloc() { 119 int x; 120 #pragma omp target teams firstprivate(x) map(alloc:x) 121 ; 122 } 123 124 // HOST: define {{.*}}mapArray 125 // HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R0]].region_id{{.*}} @[[MAPTYPES_ARRAY_R0]] 126 // HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R1]].region_id{{.*}} @[[MAPTYPES_ARRAY_R1]] 127 // 128 // CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R0]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) 129 // CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R0:.omp_outlined.[.0-9]*]] 130 // 131 // CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) 132 // CHECK: %.omp.reduction.red_list = alloca [1 x i8*] 133 // 134 // CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R1]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) 135 // CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R1:.omp_outlined.[.0-9]*]] 136 // 137 // CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z) 138 // CHECK: %.omp.reduction.red_list = alloca [1 x i8*] 139 void mapArray() { 140 int x[77], y[88], z[99]; 141 #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z) 142 ; 143 #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(to:x,y,z) 144 ; 145 } 146 147 # if HAS_INT128 148 // HOST-INT128: define {{.*}}mapInt128 149 // HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R0]].region_id{{.*}} @[[MAPTYPES_INT128_R0]] 150 // HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R1]].region_id{{.*}} @[[MAPTYPES_INT128_R1]] 151 // 152 // INT128: define {{.*}} void @[[OFFLOAD_INT128_R0]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) 153 // INT128: call void ({{.*}}@[[OUTLINE_INT128_R0:.omp_outlined.[.0-9]*]] 154 // 155 // INT128: define {{.*}} void @[[OUTLINE_INT128_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) 156 // INT128: %.omp.reduction.red_list = alloca [1 x i8*] 157 // 158 // INT128: define {{.*}} void @[[OFFLOAD_INT128_R1]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) 159 // INT128: call void ({{.*}}@[[OUTLINE_INT128_R1:.omp_outlined.[.0-9]*]] 160 // 161 // INT128: define {{.*}} void @[[OUTLINE_INT128_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z) 162 // INT128: %.omp.reduction.red_list = alloca [1 x i8*] 163 void mapInt128() { 164 __int128 x, y, z; 165 #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z) 166 ; 167 #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(from:x,y,z) 168 ; 169 } 170 # endif 171 #endif 172