clang/test/OpenMP/nvptx_allocate_messages.cpp

   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-target-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-target-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
  15 int 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
  57 template <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
  74 int 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