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 -DDEVICE -fopenmp -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -fsyntax-only %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc 3 // RUN: %clang_cc1 -verify -DDEVICE -DREQUIRES -fopenmp -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -fsyntax-only %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc 4 #if !defined(DEVICE) || defined(REQUIRES) 5 // expected-no-diagnostics 6 #endif // DEVICE 7 8 #ifndef HEADER 9 #define HEADER 10 11 #if defined(REQUIRES) && defined(DEVICE) 12 #pragma omp requires dynamic_allocators 13 #endif // REQUIRES && DEVICE 14 bar()15int bar() { 16 int res = 0; 17 #if defined(DEVICE) && !defined(REQUIRES) 18 // expected-error@+2 {{expected an 'allocator' clause inside of the target region; provide an 'allocator' clause or use 'requires' directive with the 'dynamic_allocators' clause}} 19 #endif // DEVICE && !REQUIRES 20 #pragma omp allocate(res) 21 return 0; 22 } 23 24 #pragma omp declare target 25 typedef void **omp_allocator_handle_t; 26 extern const omp_allocator_handle_t omp_null_allocator; 27 extern const omp_allocator_handle_t omp_default_mem_alloc; 28 extern const omp_allocator_handle_t omp_large_cap_mem_alloc; 29 extern const omp_allocator_handle_t omp_const_mem_alloc; 30 extern const omp_allocator_handle_t omp_high_bw_mem_alloc; 31 extern const omp_allocator_handle_t omp_low_lat_mem_alloc; 32 extern const omp_allocator_handle_t omp_cgroup_mem_alloc; 33 extern const omp_allocator_handle_t omp_pteam_mem_alloc; 34 extern const omp_allocator_handle_t omp_thread_mem_alloc; 35 36 struct St{ 37 int a; 38 }; 39 40 struct St1{ 41 int a; 42 static int b; 43 #pragma omp allocate(b) allocator(omp_default_mem_alloc) 44 } d; 45 46 int a, b, c; 47 #pragma omp allocate(a) allocator(omp_large_cap_mem_alloc) 48 #pragma omp allocate(b) allocator(omp_const_mem_alloc) 49 #pragma omp allocate(d, c) allocator(omp_high_bw_mem_alloc) 50 51 template <class T> 52 struct ST { 53 static T m; 54 #pragma omp allocate(m) allocator(omp_low_lat_mem_alloc) 55 }; 56 foo()57template <class T> T foo() { 58 T v; 59 #pragma omp allocate(v) allocator(omp_cgroup_mem_alloc) 60 v = ST<T>::m; 61 #if defined(DEVICE) && !defined(REQUIRES) 62 // expected-error@+2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}} 63 #endif // DEVICE && !REQUIRES 64 #pragma omp parallel private(v) allocate(v) 65 v = 0; 66 return v; 67 } 68 69 namespace ns{ 70 int a; 71 } 72 #pragma omp allocate(ns::a) allocator(omp_pteam_mem_alloc) 73 main()74int main () { 75 static int a; 76 #pragma omp allocate(a) allocator(omp_thread_mem_alloc) 77 a=2; 78 double b = 3; 79 #if defined(DEVICE) && !defined(REQUIRES) 80 // expected-error@+2 {{expected an 'allocator' clause inside of the target region; provide an 'allocator' clause or use 'requires' directive with the 'dynamic_allocators' clause}} 81 #endif // DEVICE && !REQUIRES 82 #pragma omp allocate(b) 83 #if defined(DEVICE) && !defined(REQUIRES) 84 // expected-note@+2 2{{called by 'main'}} 85 #endif // DEVICE && !REQUIRES 86 return (foo<int>() + bar()); 87 } 88 89 extern template int ST<int>::m; 90 #pragma omp end declare target 91 #endif 92