1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
| // Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s
// expected-no-diagnostics
// CHECK-DAG: [[T:%.+]] = type {{.+}}, {{fp128|ppc_fp128}},
// CHECK-DAG: [[T1:%.+]] = type {{.+}}, i128, i128,
#ifndef _ARCH_PPC
typedef __float128 BIGTYPE;
#else
typedef long double BIGTYPE;
#endif
struct T {
char a;
BIGTYPE f;
char c;
T() : a(12), f(15) {}
T &operator+(T &b) { f += b.a; return *this;}
};
struct T1 {
char a;
__int128 f;
__int128 f1;
char c;
T1() : a(12), f(15) {}
T1 &operator+(T1 &b) { f += b.a; return *this;}
};
#pragma omp declare target
T a = T();
T f = a;
// CHECK: define{{ dso_local | }}void @{{.+}}foo{{.+}}([[T]]* byval([[T]]) align {{.+}})
void foo(T a = T()) {
return;
}
// CHECK: define{{ dso_local | }}[6 x i64] @{{.+}}bar{{.+}}()
T bar() {
// CHECK: bitcast [[T]]* %{{.+}} to [6 x i64]*
// CHECK-NEXT: load [6 x i64], [6 x i64]* %{{.+}},
// CHECK-NEXT: ret [6 x i64]
return T();
}
// CHECK: define{{ dso_local | }}void @{{.+}}baz{{.+}}()
void baz() {
// CHECK: call [6 x i64] @{{.+}}bar{{.+}}()
// CHECK-NEXT: bitcast [[T]]* %{{.+}} to [6 x i64]*
// CHECK-NEXT: store [6 x i64] %{{.+}}, [6 x i64]* %{{.+}},
T t = bar();
}
T1 a1 = T1();
T1 f1 = a1;
// CHECK: define{{ dso_local | }}void @{{.+}}foo1{{.+}}([[T1]]* byval([[T1]]) align {{.+}})
void foo1(T1 a = T1()) {
return;
}
// CHECK: define{{ dso_local | }}[[T1]] @{{.+}}bar1{{.+}}()
T1 bar1() {
// CHECK: load [[T1]], [[T1]]*
// CHECK-NEXT: ret [[T1]]
return T1();
}
// CHECK: define{{ dso_local | }}void @{{.+}}baz1{{.+}}()
void baz1() {
// CHECK: call [[T1]] @{{.+}}bar1{{.+}}()
T1 t = bar1();
}
#pragma omp end declare target
BIGTYPE foo(BIGTYPE f) {
#pragma omp target map(f)
f = 1;
return f;
}
// CHECK: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l75([[BIGTYPE:.+]]*
// CHECK: store [[BIGTYPE]] {{0xL00000000000000003FFF000000000000|0xM3FF00000000000000000000000000000}}, [[BIGTYPE]]* %
|