235 lines · cpp
1// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP452// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping3// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP454// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -DOMP5 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP505// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -DOMP5 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping6// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -DOMP5 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP507// RUN: %clang_cc1 -verify -fopenmp -DOMP51 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP518// RUN: %clang_cc1 -fopenmp -DOMP51 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping9// RUN: %clang_cc1 -fopenmp -DOMP51 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP5110 11// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP4512// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping13// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP4514// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -DOMP5 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP5015// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -DOMP5 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping16// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -DOMP5 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP5017// RUN: %clang_cc1 -verify -fopenmp-simd -DOMP51 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP5118// RUN: %clang_cc1 -fopenmp-simd -DOMP51 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping19// RUN: %clang_cc1 -fopenmp-simd -DOMP51 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP5120// expected-no-diagnostics21 22#ifndef HEADER23#define HEADER24 25typedef 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 36void foo() {}37 38struct S {39 S(): a(0) {}40 S(int v) : a(v) {}41 int a;42 typedef int type;43};44 45template <typename T>46class S7 : public T {47protected:48 T a;49 S7() : a(0) {}50 51public:52 S7(typename T::type v) : a(v) {53#pragma omp target teams distribute simd private(a) private(this->a) private(T::a)54 for (int k = 0; k < a.a; ++k)55 ++this->a.a;56 }57 S7 &operator=(S7 &s) {58 int k;59#pragma omp target teams distribute simd allocate(a) private(a) private(this->a) linear(k) allocate(k)60 for (k = 0; k < s.a.a; ++k)61 ++s.a.a;62 return *this;63 }64 void foo() {65 int b, argv, d, c, e, f;66#pragma omp target teams distribute simd private(b), firstprivate(argv) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)67 for (int k = 0; k < a.a; ++k)68 ++a.a;69 }70 void bar() {71 int arr[10];72 const int alen = 16;73 const int slen1 = 8;74 const int slen2 = 8;75#pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr:alen)76 for (int k = 0; k < a.a; ++k)77 ++a.a;78 }79};80// CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a) private(T::a)81// CHECK: #pragma omp target teams distribute simd allocate(this->a) private(this->a) private(this->a) linear(k) allocate(k)82// CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)83// CHECK: #pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr: alen)84// CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a) private(this->S::a)85 86class S8 : public S7<S> {87 S8() {}88 89public:90 S8(int v) : S7<S>(v){91#pragma omp target teams distribute simd private(a) private(this->a) private(S7<S>::a)92 for (int k = 0; k < a.a; ++k)93 ++this->a.a;94 }95 S8 &operator=(S8 &s) {96#pragma omp target teams distribute simd private(a) private(this->a)97 for (int k = 0; k < s.a.a; ++k)98 ++s.a.a;99 return *this;100 }101 void bar() {102 int b, argv, d, c, e, f;103#pragma omp target teams distribute simd private(b), firstprivate(argv) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)104 for (int k = 0; k < a.a; ++k)105 ++a.a;106 }107 void foo() {108 const int alen = 16;109 const int slen1 = 8;110 const int slen2 = 8;111 int arr[10];112#pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr:alen)113 for (int k = 0; k < a.a; ++k)114 ++a.a;115 }116};117// CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a) private(this->S7<S>::a)118// CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a)119// CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)120// CHECK: #pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr: alen)121 122template <class T, int N>123T tmain(T argc) {124 T b = argc, c, d, e, f, g;125 static T a;126// CHECK: static T a;127 const T clen = 5;128 const T alen = 16;129 int arr[10];130#pragma omp target teams distribute simd131 for (int i=0; i < 2; ++i)132 a = 2;133// CHECK: #pragma omp target teams distribute simd{{$}}134// CHECK-NEXT: for (int i = 0; i < 2; ++i)135// CHECK-NEXT: a = 2;136#pragma omp target teams distribute simd private(argc, b), firstprivate(c, d), collapse(2)137 for (int i = 0; i < 10; ++i)138 for (int j = 0; j < 10; ++j)139 foo();140// CHECK: #pragma omp target teams distribute simd private(argc,b) firstprivate(c,d) collapse(2)141// CHECK-NEXT: for (int i = 0; i < 10; ++i)142// CHECK-NEXT: for (int j = 0; j < 10; ++j)143// CHECK-NEXT: foo();144 for (int i = 0; i < 10; ++i)145 foo();146// CHECK: for (int i = 0; i < 10; ++i)147// CHECK-NEXT: foo();148#pragma omp target teams distribute simd149 for (int i = 0; i < 10; ++i)150 foo();151// CHECK: #pragma omp target teams distribute simd152// CHECK-NEXT: for (int i = 0; i < 10; ++i)153// CHECK-NEXT: foo();154#pragma omp target teams distribute simd private(b), firstprivate(argc) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)155 for (int k = 0; k < 10; ++k)156 e += d + argc;157// CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argc) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)158// CHECK-NEXT: for (int k = 0; k < 10; ++k)159// CHECK-NEXT: e += d + argc;160#pragma omp target teams distribute simd simdlen(clen-1)161 for (int k = 0; k < 10; ++k)162 e += d + argc;163// CHECK: #pragma omp target teams distribute simd simdlen(clen - 1)164// CHECK-NEXT: for (int k = 0; k < 10; ++k)165// CHECK-NEXT: e += d + argc;166#pragma omp target teams distribute simd safelen(clen-1) aligned(arr:alen)167 for (int k = 0; k < 10; ++k)168 e += d + argc + arr[k];169// CHECK: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: alen)170// CHECK-NEXT: for (int k = 0; k < 10; ++k)171// CHECK-NEXT: e += d + argc + arr[k];172 return T();173}174 175int main (int argc, char **argv) {176 int b = argc, c, d, e, f, g;177 static int a;178// CHECK: static int a;179 const int clen = 5;180 const int N = 10;181 int arr[10];182#pragma omp target teams distribute simd183 for (int i=0; i < 2; ++i)184 a = 2;185// CHECK: #pragma omp target teams distribute simd186// CHECK-NEXT: for (int i = 0; i < 2; ++i)187// CHECK-NEXT: a = 2;188#pragma omp target teams distribute simd private(argc,b),firstprivate(argv, c), collapse(2)189 for (int i = 0; i < 10; ++i)190 for (int j = 0; j < 10; ++j)191 foo();192// CHECK: #pragma omp target teams distribute simd private(argc,b) firstprivate(argv,c) collapse(2)193// CHECK-NEXT: for (int i = 0; i < 10; ++i)194// CHECK-NEXT: for (int j = 0; j < 10; ++j)195// CHECK-NEXT: foo();196 for (int i = 0; i < 10; ++i)197 foo();198// CHECK: for (int i = 0; i < 10; ++i)199// CHECK-NEXT: foo();200#pragma omp target teams distribute simd201 for (int i = 0; i < 10; ++i)foo();202// CHECK: #pragma omp target teams distribute simd203// CHECK-NEXT: for (int i = 0; i < 10; ++i)204// CHECK-NEXT: foo();205#pragma omp target teams distribute simd private(b), firstprivate(argc) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)206 for (int k = 0; k < 10; ++k)207 e += d + argc;208// CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argc) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)209// CHECK-NEXT: for (int k = 0; k < 10; ++k)210// CHECK-NEXT: e += d + argc;211#pragma omp target teams distribute simd simdlen(clen-1)212 for (int k = 0; k < 10; ++k)213 e += d + argc;214// CHECK: #pragma omp target teams distribute simd simdlen(clen - 1)215// CHECK-NEXT: for (int k = 0; k < 10; ++k)216// CHECK-NEXT: e += d + argc;217#ifdef OMP51218#pragma omp target teams distribute simd safelen(clen-1) aligned(arr:N+6) if(simd:argc) nontemporal(argc, c, d) order(unconstrained:concurrent) allocate(omp_low_lat_mem_alloc:e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)219#elif OMP5220#pragma omp target teams distribute simd safelen(clen-1) aligned(arr:N+6) if(simd:argc) nontemporal(argc, c, d) order(concurrent) allocate(omp_low_lat_mem_alloc:e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)221#else222#pragma omp target teams distribute simd safelen(clen-1) aligned(arr:N+6)223#endif // OMP51224 for (int k = 0; k < 10; ++k)225 e += d + argc + arr[k];226// OMP45: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: N + 6)227// OMP50: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: N + 6) if(simd: argc) nontemporal(argc,c,d) order(concurrent) allocate(omp_low_lat_mem_alloc: e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)228// OMP51: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: N + 6) if(simd: argc) nontemporal(argc,c,d) order(unconstrained: concurrent) allocate(omp_low_lat_mem_alloc: e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)229// CHECK-NEXT: for (int k = 0; k < 10; ++k)230// CHECK-NEXT: e += d + argc + arr[k];231 return (0);232}233 234#endif235