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_null_allocator;
11 extern const omp_allocator_handle_t omp_default_mem_alloc;
12 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
13 extern const omp_allocator_handle_t omp_const_mem_alloc;
14 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
15 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
16 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
17 extern const omp_allocator_handle_t omp_pteam_mem_alloc;
18 extern const omp_allocator_handle_t omp_thread_mem_alloc;
19
20 // CHECK-DAG: @{{.+}}St1{{.+}}b{{.+}} = external global i32,
21 // CHECK-DAG: @a ={{ hidden | }}global i32 0,
22 // CHECK-DAG: @b ={{ hidden | }}addrspace(4) global i32 0,
23 // CHECK-DAG: @c ={{ hidden | }}global i32 0,
24 // CHECK-DAG: @d ={{ hidden | }}global %struct.St1 zeroinitializer,
25 // CHECK-DAG: @{{.+}}ns{{.+}}a{{.+}} ={{ hidden | }}addrspace(3) global i32 0,
26 // CHECK-DAG: @{{.+}}main{{.+}}a{{.*}} = internal global i32 0,
27 // CHECK-DAG: @{{.+}}ST{{.+}}m{{.+}} = external global i32,
28 // CHECK-DAG: @bar_c = internal global i32 0,
29 // CHECK-DAG: @bar_b = internal addrspace(3) global double 0.000000e+00,
30 struct St{
31 int a;
32 };
33
34 struct St1{
35 int a;
36 static int b;
37 #pragma omp allocate(b) allocator(omp_default_mem_alloc)
38 } d;
39
40 int a, b, c;
41 #pragma omp allocate(a) allocator(omp_large_cap_mem_alloc)
42 #pragma omp allocate(b) allocator(omp_const_mem_alloc)
43 #pragma omp allocate(d, c) allocator(omp_high_bw_mem_alloc)
44
45 template <class T>
46 struct ST {
47 static T m;
48 #pragma omp allocate(m) allocator(omp_low_lat_mem_alloc)
49 };
50
foo()51 template <class T> T foo() {
52 T v;
53 #pragma omp allocate(v) allocator(omp_cgroup_mem_alloc)
54 v = ST<T>::m;
55 return v;
56 }
57
58 namespace ns{
59 int a;
60 }
61 #pragma omp allocate(ns::a) allocator(omp_pteam_mem_alloc)
62
63 // CHECK-LABEL: @main
main()64 int main () {
65 // CHECK: alloca double,
66 static int a;
67 #pragma omp allocate(a) allocator(omp_thread_mem_alloc)
68 a=2;
69 double b = 3;
70 float c;
71 #pragma omp allocate(b) allocator(omp_default_mem_alloc)
72 #pragma omp allocate(c) allocator(omp_cgroup_mem_alloc)
73 return (foo<int>());
74 }
75
76 // CHECK: define {{.*}}i32 @{{.+}}foo{{.+}}()
77 // CHECK-NOT: alloca i32,
78
79 extern template int ST<int>::m;
80
81 void baz(float &);
82
83 // CHECK: define{{ hidden | }}void @{{.+}}bar{{.+}}()
bar()84 void bar() {
85 // CHECK: alloca float,
86 float bar_a;
87 // CHECK: alloca double,
88 double bar_b;
89 int bar_c;
90 #pragma omp allocate(bar_c) allocator(omp_cgroup_mem_alloc)
91 // CHECK: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}})
92 #pragma omp parallel private(bar_a, bar_b) allocate(omp_thread_mem_alloc \
93 : bar_a) allocate(omp_pteam_mem_alloc \
94 : bar_b)
95 {
96 bar_b = bar_a;
97 baz(bar_a);
98 }
99 // CHECK: define internal void [[OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}})
100 // CHECK-NOT: alloca double,
101 // CHECK: alloca float,
102 // CHECK-NOT: alloca double,
103 // CHECK: load float, float* %
104 // CHECK: store double {{.+}}, double* addrspacecast (double addrspace(3)* @bar_b to double*),
105 }
106
107 #pragma omp end declare target
108 #endif
109