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 3, i64 3]
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 3]
30 // HOST: @[[MAPTYPES_ARRAY_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 33, i64 33, i64 1]
31 // HOST-INT128: @[[MAPTYPES_INT128_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 35, i64 35, i64 3]
32 // HOST-INT128: @[[MAPTYPES_INT128_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [3 x i64] [i64 34, i64 34, i64 2]
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.)
mapWithPrivate()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)
mapWithFirstprivate()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*]
mapWithReduction()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)
mapFrom()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)
mapTo()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)
mapAlloc()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*]
mapArray()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*]
mapInt128()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