1 // RUN: %clang_cc1 -verify -fopenmp -triple x86_64-apple-darwin10.6.0 -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc -o %t-host.bc %s
2 // RUN: %clang_cc1 -verify -fopenmp -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - -disable-llvm-optzns | FileCheck %s
3 // expected-no-diagnostics
4
5 #ifndef HEADER
6 #define HEADER
7
8 #pragma omp declare target
9 typedef void **omp_allocator_handle_t;
10 extern const omp_allocator_handle_t omp_default_mem_alloc;
11 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
12 extern const omp_allocator_handle_t omp_const_mem_alloc;
13 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
14 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
15 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
16 extern const omp_allocator_handle_t omp_pteam_mem_alloc;
17 extern const omp_allocator_handle_t omp_thread_mem_alloc;
18
19 // CHECK-DAG: @{{.+}}St1{{.+}}b{{.+}} = external global i32,
20 // CHECK-DAG: @a ={{ dso_local | }}global i32 0,
21 // CHECK-DAG: @b ={{ dso_local | }}addrspace(4) global i32 0,
22 // CHECK-DAG: @c ={{ dso_local | }}global i32 0,
23 // CHECK-DAG: @d ={{ dso_local | }}global %struct.St1 zeroinitializer,
24 // CHECK-DAG: @{{.+}}ns{{.+}}a{{.+}} ={{ dso_local | }}addrspace(3) global i32 0,
25 // CHECK-DAG: @{{.+}}main{{.+}}a{{.*}} = internal global i32 0,
26 // CHECK-DAG: @{{.+}}ST{{.+}}m{{.+}} = external global i32,
27 // CHECK-DAG: @bar_c = internal global i32 0,
28 // CHECK-DAG: @bar_b = internal addrspace(3) global double 0.000000e+00,
29 struct St{
30 int a;
31 };
32
33 struct St1{
34 int a;
35 static int b;
36 #pragma omp allocate(b) allocator(omp_default_mem_alloc)
37 } d;
38
39 int a, b, c;
40 #pragma omp allocate(a) allocator(omp_large_cap_mem_alloc)
41 #pragma omp allocate(b) allocator(omp_const_mem_alloc)
42 #pragma omp allocate(d, c) allocator(omp_high_bw_mem_alloc)
43
44 template <class T>
45 struct ST {
46 static T m;
47 #pragma omp allocate(m) allocator(omp_low_lat_mem_alloc)
48 };
49
foo()50 template <class T> T foo() {
51 T v;
52 #pragma omp allocate(v) allocator(omp_cgroup_mem_alloc)
53 v = ST<T>::m;
54 return v;
55 }
56
57 namespace ns{
58 int a;
59 }
60 #pragma omp allocate(ns::a) allocator(omp_pteam_mem_alloc)
61
62 // CHECK-LABEL: @main
main()63 int main () {
64 // CHECK: alloca double,
65 static int a;
66 #pragma omp allocate(a) allocator(omp_thread_mem_alloc)
67 a=2;
68 double b = 3;
69 float c;
70 #pragma omp allocate(b) allocator(omp_default_mem_alloc)
71 #pragma omp allocate(c) allocator(omp_cgroup_mem_alloc)
72 return (foo<int>());
73 }
74
75 // CHECK: define {{.*}}i32 @{{.+}}foo{{.+}}()
76 // CHECK-NOT: alloca i32,
77
78 extern template int ST<int>::m;
79
80 void baz(float &);
81
82 // CHECK: define{{ dso_local | }}void @{{.+}}bar{{.+}}()
bar()83 void bar() {
84 // CHECK: alloca float,
85 float bar_a;
86 // CHECK: alloca double,
87 double bar_b;
88 int bar_c;
89 #pragma omp allocate(bar_c) allocator(omp_cgroup_mem_alloc)
90 // CHECK: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}})
91 #pragma omp parallel private(bar_a, bar_b) allocate(omp_thread_mem_alloc \
92 : bar_a) allocate(omp_pteam_mem_alloc \
93 : bar_b)
94 {
95 bar_b = bar_a;
96 baz(bar_a);
97 }
98 // CHECK: define internal void [[OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}})
99 // CHECK-NOT: alloca double,
100 // CHECK: alloca float,
101 // CHECK-NOT: alloca double,
102 // CHECK: load float, float* %
103 // CHECK: store double {{.+}}, double addrspace(3)* @bar_b,
104 }
105
106 #pragma omp end declare target
107 #endif
108