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