92 lines · 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 %s2// 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.bc3// 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.bc4#if !defined(DEVICE) || defined(REQUIRES)5// expected-no-diagnostics6#endif // DEVICE7 8#ifndef HEADER9#define HEADER10 11#if defined(REQUIRES) && defined(DEVICE)12#pragma omp requires dynamic_allocators13#endif // REQUIRES && DEVICE14 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 && !REQUIRES20#pragma omp allocate(res)21 return 0;22}23 24#pragma omp declare target25typedef void **omp_allocator_handle_t;26extern const omp_allocator_handle_t omp_null_allocator;27extern const omp_allocator_handle_t omp_default_mem_alloc;28extern const omp_allocator_handle_t omp_large_cap_mem_alloc;29extern const omp_allocator_handle_t omp_const_mem_alloc;30extern const omp_allocator_handle_t omp_high_bw_mem_alloc;31extern const omp_allocator_handle_t omp_low_lat_mem_alloc;32extern const omp_allocator_handle_t omp_cgroup_mem_alloc;33extern const omp_allocator_handle_t omp_pteam_mem_alloc;34extern const omp_allocator_handle_t omp_thread_mem_alloc;35 36struct St{37 int a;38};39 40struct St1{41 int a;42 static int b;43#pragma omp allocate(b) allocator(omp_default_mem_alloc)44} d;45 46int 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 51template <class T>52struct ST {53 static T m;54 #pragma omp allocate(m) allocator(omp_low_lat_mem_alloc)55};56 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 && !REQUIRES64#pragma omp parallel private(v) allocate(v)65 v = 0;66 return v;67}68 69namespace ns{70 int a;71}72#pragma omp allocate(ns::a) allocator(omp_pteam_mem_alloc)73 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 && !REQUIRES82#pragma omp allocate(b)83#if defined(DEVICE) && !defined(REQUIRES)84// expected-note@+2 2{{called by 'main'}}85#endif // DEVICE && !REQUIRES86 return (foo<int>() + bar());87}88 89extern template int ST<int>::m;90#pragma omp end declare target91#endif92