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 
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
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{{.+}}()
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