339 lines · cpp
1// RUN: %clang_cc1 -verify -fopenmp -std=c++11 \2// RUN: -ast-print %s | FileCheck %s3 4// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 \5// RUN: -emit-pch -o %t %s6 7// RUN: %clang_cc1 -fopenmp -std=c++11 \8// RUN: -include-pch %t -verify %s -ast-print | FileCheck %s9 10// RUN: %clang_cc1 -verify -fopenmp-simd \11// RUN: -std=c++11 -ast-print %s | FileCheck %s12 13// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 \14// RUN: -emit-pch -o %t %s15 16// RUN: %clang_cc1 -fopenmp-simd -std=c++11 \17// RUN: -include-pch %t -verify %s -ast-print | FileCheck %s18 19// expected-no-diagnostics20 21#ifndef HEADER22#define HEADER23 24struct ST {25 int *a;26};27typedef int arr[10];28typedef ST STarr[10];29struct SA {30 const int da[5] = { 0 };31 ST g[10];32 STarr &rg = g;33 int i;34 int &j = i;35 int *k = &j;36 int *&z = k;37 int aa[10];38 arr &raa = aa;39 void func(int arg) {40#pragma omp target has_device_addr(k)41 {}42#pragma omp target has_device_addr(z)43 {}44#pragma omp target has_device_addr(aa) // OK45 {}46#pragma omp target has_device_addr(raa) // OK47 {}48#pragma omp target has_device_addr(g) // OK49 {}50#pragma omp target has_device_addr(rg) // OK51 {}52#pragma omp target has_device_addr(da) // OK53 {}54 return;55 }56};57// CHECK: struct SA58// CHECK-NEXT: const int da[5] = {0};59// CHECK-NEXT: ST g[10];60// CHECK-NEXT: STarr &rg = this->g;61// CHECK-NEXT: int i;62// CHECK-NEXT: int &j = this->i;63// CHECK-NEXT: int *k = &this->j;64// CHECK-NEXT: int *&z = this->k;65// CHECK-NEXT: int aa[10];66// CHECK-NEXT: arr &raa = this->aa;67// CHECK-NEXT: func(68// CHECK-NEXT: #pragma omp target has_device_addr(this->k)69// CHECK-NEXT: {70// CHECK-NEXT: }71// CHECK-NEXT: #pragma omp target has_device_addr(this->z)72// CHECK-NEXT: {73// CHECK-NEXT: }74// CHECK-NEXT: #pragma omp target has_device_addr(this->aa)75// CHECK-NEXT: {76// CHECK-NEXT: }77// CHECK-NEXT: #pragma omp target has_device_addr(this->raa)78// CHECK-NEXT: {79// CHECK-NEXT: }80// CHECK-NEXT: #pragma omp target has_device_addr(this->g)81// CHECK-NEXT: {82// CHECK-NEXT: }83// CHECK-NEXT: #pragma omp target has_device_addr(this->rg)84// CHECK-NEXT: {85// CHECK-NEXT: }86// CHECK-NEXT: #pragma omp target has_device_addr(this->da)87 88struct SB {89 unsigned A;90 unsigned B;91 float Arr[100];92 float *Ptr;93 float *foo() {94 return &Arr[0];95 }96};97 98struct SC {99 unsigned A : 2;100 unsigned B : 3;101 unsigned C;102 unsigned D;103 float Arr[100];104 SB S;105 SB ArrS[100];106 SB *PtrS;107 SB *&RPtrS;108 float *Ptr;109 110 SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}111};112 113union SD {114 unsigned A;115 float B;116};117 118struct S1;119extern S1 a;120class S2 {121 mutable int a;122public:123 S2():a(0) { }124 S2(S2 &s2):a(s2.a) { }125 static float S2s;126 static const float S2sc;127};128const float S2::S2sc = 0;129const S2 b;130const S2 ba[5];131class S3 {132 int a;133public:134 S3():a(0) { }135 S3(S3 &s3):a(s3.a) { }136};137const S3 c;138const S3 ca[5];139extern const int f;140class S4 {141 int a;142 S4();143 S4(const S4 &s4);144public:145 S4(int v):a(v) { }146};147class S5 {148 int a;149 S5():a(0) {}150 S5(const S5 &s5):a(s5.a) { }151public:152 S5(int v):a(v) { }153};154 155S3 h;156#pragma omp threadprivate(h)157 158typedef struct {159 int a;160} S6;161 162template <typename T>163T tmain(T argc) {164 const T da[5] = { 0 };165 S6 h[10];166 auto &rh = h;167 T i;168 T &j = i;169 T *k = &j;170 T *&z = k;171 T aa[10];172 auto &raa = aa;173#pragma omp target has_device_addr(k)174 {}175#pragma omp target has_device_addr(z)176 {}177#pragma omp target has_device_addr(aa)178 {}179#pragma omp target has_device_addr(raa)180 {}181#pragma omp target has_device_addr(h)182 {}183#pragma omp target has_device_addr(rh)184 {}185#pragma omp target has_device_addr(da)186 {}187 return 0;188}189 190// CHECK: template<> int tmain<int>(int argc) {191// CHECK-NEXT: const int da[5] = {0};192// CHECK-NEXT: S6 h[10];193// CHECK-NEXT: auto &rh = h;194// CHECK-NEXT: int i;195// CHECK-NEXT: int &j = i;196// CHECK-NEXT: int *k = &j;197// CHECK-NEXT: int *&z = k;198// CHECK-NEXT: int aa[10];199// CHECK-NEXT: auto &raa = aa;200// CHECK-NEXT: #pragma omp target has_device_addr(k)201// CHECK-NEXT: {202// CHECK-NEXT: }203// CHECK-NEXT: #pragma omp target has_device_addr(z)204// CHECK-NEXT: {205// CHECK-NEXT: }206// CHECK-NEXT: #pragma omp target has_device_addr(aa)207// CHECK-NEXT: {208// CHECK-NEXT: }209// CHECK-NEXT: #pragma omp target has_device_addr(raa)210// CHECK-NEXT: {211// CHECK-NEXT: }212// CHECK-NEXT: #pragma omp target has_device_addr(h)213// CHECK-NEXT: {214// CHECK-NEXT: }215// CHECK-NEXT: #pragma omp target has_device_addr(rh)216// CHECK-NEXT: {217// CHECK-NEXT: }218// CHECK-NEXT: #pragma omp target has_device_addr(da)219 220// CHECK: template<> int *tmain<int *>(int *argc) {221// CHECK-NEXT: int *const da[5] = {0};222// CHECK-NEXT: S6 h[10];223// CHECK-NEXT: auto &rh = h;224// CHECK-NEXT: int *i;225// CHECK-NEXT: int *&j = i;226// CHECK-NEXT: int **k = &j;227// CHECK-NEXT: int **&z = k;228// CHECK-NEXT: int *aa[10];229// CHECK-NEXT: auto &raa = aa;230// CHECK-NEXT: #pragma omp target has_device_addr(k)231// CHECK-NEXT: {232// CHECK-NEXT: }233// CHECK-NEXT: #pragma omp target has_device_addr(z)234// CHECK-NEXT: {235// CHECK-NEXT: }236// CHECK-NEXT: #pragma omp target has_device_addr(aa)237// CHECK-NEXT: {238// CHECK-NEXT: }239// CHECK-NEXT: #pragma omp target has_device_addr(raa)240// CHECK-NEXT: {241// CHECK-NEXT: }242// CHECK-NEXT: #pragma omp target has_device_addr(h)243// CHECK-NEXT: {244// CHECK-NEXT: }245// CHECK-NEXT: #pragma omp target has_device_addr(rh)246// CHECK-NEXT: {247// CHECK-NEXT: }248// CHECK-NEXT: #pragma omp target has_device_addr(da)249 250// CHECK-LABEL: int main(int argc, char **argv) {251int main(int argc, char **argv) {252 const int da[5] = { 0 };253 S6 h[10];254 auto &rh = h;255 int i;256 int &j = i;257 int *k = &j;258 int *&z = k;259 int aa[10];260 auto &raa = aa;261// CHECK-NEXT: const int da[5] = {0};262// CHECK-NEXT: S6 h[10];263// CHECK-NEXT: auto &rh = h;264// CHECK-NEXT: int i;265// CHECK-NEXT: int &j = i;266// CHECK-NEXT: int *k = &j;267// CHECK-NEXT: int *&z = k;268// CHECK-NEXT: int aa[10];269// CHECK-NEXT: auto &raa = aa;270#pragma omp target has_device_addr(k)271// CHECK-NEXT: #pragma omp target has_device_addr(k)272 {}273// CHECK-NEXT: {274// CHECK-NEXT: }275#pragma omp target has_device_addr(z)276// CHECK-NEXT: #pragma omp target has_device_addr(z)277 {}278// CHECK-NEXT: {279// CHECK-NEXT: }280#pragma omp target has_device_addr(aa)281// CHECK-NEXT: #pragma omp target has_device_addr(aa)282 {}283// CHECK-NEXT: {284// CHECK-NEXT: }285#pragma omp target has_device_addr(raa)286// CHECK-NEXT: #pragma omp target has_device_addr(raa)287 {}288// CHECK-NEXT: {289// CHECK-NEXT: }290#pragma omp target has_device_addr(h)291// CHECK-NEXT: #pragma omp target has_device_addr(h)292 {}293// CHECK-NEXT: {294// CHECK-NEXT: }295#pragma omp target has_device_addr(rh)296// CHECK-NEXT: #pragma omp target has_device_addr(rh)297 {}298// CHECK-NEXT: {299// CHECK-NEXT: }300#pragma omp target has_device_addr(da)301// CHECK-NEXT: #pragma omp target has_device_addr(da)302 {}303// CHECK-NEXT: {304// CHECK-NEXT: }305// CHECK-NEXT: #pragma omp target has_device_addr(da[1:3])306// CHECK-NEXT: {307// CHECK-NEXT: }308#pragma omp target has_device_addr(da[1:3])309 {}310 return tmain<int>(argc) + *tmain<int *>(&argc);311}312 313struct SomeKernel {314 int targetDev;315 float devPtr;316 SomeKernel();317 ~SomeKernel();318 319 template<unsigned int nRHS>320 void apply() {321 #pragma omp target has_device_addr(devPtr) device(targetDev)322 {323 }324// CHECK: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)325// CHECK-NEXT: {326// CHECK-NEXT: }327 }328// CHECK: template<> void apply<32U>() {329// CHECK: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)330// CHECK-NEXT: {331// CHECK-NEXT: }332};333 334void use_template() {335 SomeKernel aKern;336 aKern.apply<32>();337}338#endif339