ABataev updated this revision to Diff 266076.
ABataev added a comment.
Rebase + cleanup
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D80404/new/
https://reviews.llvm.org/D80404
Files:
clang/include/clang/AST/OpenMPClause.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/Sema/Sema.h
clang/lib/AST/OpenMPClause.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/Basic/OpenMPKinds.cpp
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/OpenMP/target_data_messages.c
clang/test/OpenMP/target_data_use_device_ptr_addr_ast_print.cpp
clang/test/OpenMP/target_data_use_device_ptr_addr_messages.cpp
clang/test/OpenMP/target_data_use_device_ptr_ast_print.cpp
clang/test/OpenMP/target_data_use_device_ptr_messages.cpp
clang/test/OpenMP/target_map_messages.cpp
clang/test/OpenMP/target_teams_map_messages.cpp
clang/tools/libclang/CIndex.cpp
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -198,6 +198,7 @@
__OMP_CLAUSE(exclusive, OMPExclusiveClause)
__OMP_CLAUSE(uses_allocators, OMPUsesAllocatorsClause)
__OMP_CLAUSE(affinity, OMPAffinityClause)
+__OMP_CLAUSE(use_device_addr, OMPUseDeviceAddrClause)
__OMP_CLAUSE_NO_CLASS(uniform)
__OMP_CLAUSE_NO_CLASS(device_type)
@@ -904,6 +905,7 @@
__OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, device)
__OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, map)
__OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, use_device_ptr)
+__OMP_DIRECTIVE_CLAUSE(target_data, 50, ~0, use_device_addr)
__OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, if)
__OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, device)
Index: clang/tools/libclang/CIndex.cpp
===================================================================
--- clang/tools/libclang/CIndex.cpp
+++ clang/tools/libclang/CIndex.cpp
@@ -2489,6 +2489,10 @@
const OMPUseDevicePtrClause *C) {
VisitOMPClauseList(C);
}
+void OMPClauseEnqueue::VisitOMPUseDeviceAddrClause(
+ const OMPUseDeviceAddrClause *C) {
+ VisitOMPClauseList(C);
+}
void OMPClauseEnqueue::VisitOMPIsDevicePtrClause(
const OMPIsDevicePtrClause *C) {
VisitOMPClauseList(C);
Index: clang/test/OpenMP/target_teams_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_map_messages.cpp
+++ clang/test/OpenMP/target_teams_map_messages.cpp
@@ -488,7 +488,7 @@
int y;
int to, tofrom, always;
const int (&l)[5] = da;
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
#pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
#pragma omp target data map() // expected-error {{expected expression}}
#pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
Index: clang/test/OpenMP/target_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_map_messages.cpp
+++ clang/test/OpenMP/target_map_messages.cpp
@@ -598,7 +598,7 @@
const int (&l)[5] = da;
SC1 s;
SC1 *p;
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
#pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
#pragma omp target data map() // expected-error {{expected expression}}
#pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
Index: clang/test/OpenMP/target_data_use_device_ptr_messages.cpp
===================================================================
--- clang/test/OpenMP/target_data_use_device_ptr_messages.cpp
+++ /dev/null
@@ -1,208 +0,0 @@
-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s -Wuninitialized
-
-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
-struct ST {
- int *a;
-};
-struct SA {
- const int d = 5;
- const int da[5] = { 0 };
- ST e;
- ST g[10];
- int i;
- int &j = i;
- int *k = &j;
- int *&z = k;
- int aa[10];
- void func(int arg) {
-#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
- {}
-#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
- {}
-#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
- {}
-#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
- {}
-#pragma omp target data map(i) use_device_ptr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(k) // OK
- {}
-#pragma omp target data map(i) use_device_ptr(z) // OK
- {}
-#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
- return;
- }
-};
-struct SB {
- unsigned A;
- unsigned B;
- float Arr[100];
- float *Ptr;
- float *foo() {
- return &Arr[0];
- }
-};
-
-struct SC {
- unsigned A : 2;
- unsigned B : 3;
- unsigned C;
- unsigned D;
- float Arr[100];
- SB S;
- SB ArrS[100];
- SB *PtrS;
- SB *&RPtrS;
- float *Ptr;
-
- SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
-};
-
-union SD {
- unsigned A;
- float B;
-};
-
-struct S1;
-extern S1 a;
-class S2 {
- mutable int a;
-public:
- S2():a(0) { }
- S2(S2 &s2):a(s2.a) { }
- static float S2s;
- static const float S2sc;
-};
-const float S2::S2sc = 0;
-const S2 b;
-const S2 ba[5];
-class S3 {
- int a;
-public:
- S3():a(0) { }
- S3(S3 &s3):a(s3.a) { }
-};
-const S3 c;
-const S3 ca[5];
-extern const int f;
-class S4 {
- int a;
- S4();
- S4(const S4 &s4);
-public:
- S4(int v):a(v) { }
-};
-class S5 {
- int a;
- S5():a(0) {}
- S5(const S5 &s5):a(s5.a) { }
-public:
- S5(int v):a(v) { }
-};
-
-S3 h;
-#pragma omp threadprivate(h)
-
-typedef int from;
-
-template <typename T, int I>
-T tmain(T argc) {
- const T d = 5;
- const T da[5] = { 0 };
- S4 e(4);
- S5 g(5);
- T i;
- T &j = i;
- T *k = &j;
- T *&z = k;
- T aa[10];
-#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
- {}
-#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
- {}
-#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
- {}
-#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
- {}
-#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(k) // OK
- {}
-#pragma omp target data map(i) use_device_ptr(z) // OK
- {}
-#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
- return 0;
-}
-
-int main(int argc, char **argv) {
- const int d = 5;
- const int da[5] = { 0 };
- S4 e(4);
- S5 g(5);
- int i;
- int &j = i;
- int *k = &j;
- int *&z = k;
- int aa[10];
-#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
- {}
-#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
- {}
-#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
- {}
-#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
- {}
-#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(k) // OK
- {}
-#pragma omp target data map(i) use_device_ptr(z) // OK
- {}
-#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
-#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
- {}
- return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
-}
Index: clang/test/OpenMP/target_data_use_device_ptr_addr_messages.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_data_use_device_ptr_addr_messages.cpp
@@ -0,0 +1,300 @@
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wuninitialized
+
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wuninitialized
+struct ST {
+ int *a;
+};
+struct SA {
+ const int d = 5;
+ const int da[5] = { 0 };
+ ST e;
+ ST g[10];
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+ int aa[10];
+ void func(int arg) {
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+ {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+ {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+ {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+ {}
+#pragma omp target data map(i) use_device_ptr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+ {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+ {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_addr // expected-error {{expected '(' after 'use_device_addr'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr() // expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(i) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(k) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(z) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(aa) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(e) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(g) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(k,i,j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(d) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(da) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+ return;
+ }
+};
+struct SB {
+ unsigned A;
+ unsigned B;
+ float Arr[100];
+ float *Ptr;
+ float *foo() {
+ return &Arr[0];
+ }
+};
+
+struct SC {
+ unsigned A : 2;
+ unsigned B : 3;
+ unsigned C;
+ unsigned D;
+ float Arr[100];
+ SB S;
+ SB ArrS[100];
+ SB *PtrS;
+ SB *&RPtrS;
+ float *Ptr;
+
+ SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+ unsigned A;
+ float B;
+};
+
+struct S1;
+extern S1 a;
+class S2 {
+ mutable int a;
+public:
+ S2():a(0) { }
+ S2(S2 &s2):a(s2.a) { }
+ static float S2s;
+ static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+ int a;
+public:
+ S3():a(0) { }
+ S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+ int a;
+ S4();
+ S4(const S4 &s4);
+public:
+ S4(int v):a(v) { }
+};
+class S5 {
+ int a;
+ S5():a(0) {}
+ S5(const S5 &s5):a(s5.a) { }
+public:
+ S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h)
+
+typedef int from;
+
+template <typename T, int I>
+T tmain(T argc) {
+ const T d = 5;
+ const T da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ T i;
+ T &j = i;
+ T *k = &j;
+ T *&z = k;
+ T aa[10];
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+ {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+ {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+ {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+ {}
+#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+ {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+ {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_addr // expected-error {{expected '(' after 'use_device_addr'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr() // expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(i) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(k) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(z) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(aa) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(e) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(g) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(k,i,j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(d) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(da) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+ return 0;
+}
+
+int main(int argc, char **argv) {
+ const int d = 5;
+ const int da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+ int aa[10];
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+ {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+ {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+ {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+ {}
+#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+ {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+ {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+ {}
+#pragma omp target data map(i) use_device_addr // expected-error {{expected '(' after 'use_device_addr'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr() // expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(i) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(k) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(z) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(aa) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(e) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(g) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(k,i,j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(d) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+#pragma omp target data map(i) use_device_addr(da) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+ {}
+ return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
+}
Index: clang/test/OpenMP/target_data_use_device_ptr_addr_ast_print.cpp
===================================================================
--- clang/test/OpenMP/target_data_use_device_ptr_addr_ast_print.cpp
+++ clang/test/OpenMP/target_data_use_device_ptr_addr_ast_print.cpp
@@ -1,9 +1,10 @@
-// RxUN: %clang_cc1 -verify -fopenmp -std=c++11 -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -std=c++11 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -std=c++11 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
// expected-no-diagnostics
#ifndef HEADER
@@ -16,18 +17,19 @@
int i, j;
int *k = &j;
int *&z = k;
+ int &y = i;
void func(int arg) {
-#pragma omp target data map(tofrom: i) use_device_ptr(k)
+#pragma omp target data map(tofrom: i) use_device_ptr(k) use_device_addr(i, j)
{}
-#pragma omp target data map(tofrom: i) use_device_ptr(z)
+#pragma omp target data map(tofrom: i) use_device_ptr(z) use_device_addr(k, y)
{}
return;
}
};
// CHECK: struct SA
// CHECK: void func(
-// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->k){{$}}
-// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->z)
+// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->k) use_device_addr(this->i,this->j){{$}}
+// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->z) use_device_addr(this->k,this->y)
struct SB {
unsigned A;
unsigned B;
@@ -143,13 +145,13 @@
// CHECK-NEXT: int &j = i;
// CHECK-NEXT: int *k = &j;
// CHECK-NEXT: int *&z = k;
-#pragma omp target data map(tofrom: i) use_device_ptr(k)
-// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k)
+#pragma omp target data map(tofrom: i) use_device_ptr(k) use_device_addr(i, j)
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k) use_device_addr(i,j)
{}
// CHECK-NEXT: {
// CHECK-NEXT: }
-#pragma omp target data map(tofrom: i) use_device_ptr(z)
-// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z)
+#pragma omp target data map(tofrom: i) use_device_ptr(z) use_device_addr(i, j, k[:i])
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z) use_device_addr(i,j,k[:i])
{}
return tmain<int>(argc) + (*tmain<int*>(&argc));
}
Index: clang/test/OpenMP/target_data_messages.c
===================================================================
--- clang/test/OpenMP/target_data_messages.c
+++ clang/test/OpenMP/target_data_messages.c
@@ -1,6 +1,8 @@
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - %s -Wuninitialized
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - %s -Wuninitialized
void foo() { }
@@ -13,7 +15,7 @@
int main(int argc, char **argv) {
int a;
- #pragma omp target data // expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+ #pragma omp target data // omp45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} omp50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
{}
L1:
foo();
Index: clang/lib/Serialization/ASTWriter.cpp
===================================================================
--- clang/lib/Serialization/ASTWriter.cpp
+++ clang/lib/Serialization/ASTWriter.cpp
@@ -6625,6 +6625,26 @@
}
}
+void OMPClauseWriter::VisitOMPUseDeviceAddrClause(OMPUseDeviceAddrClause *C) {
+ Record.push_back(C->varlist_size());
+ Record.push_back(C->getUniqueDeclarationsNum());
+ Record.push_back(C->getTotalComponentListNum());
+ Record.push_back(C->getTotalComponentsNum());
+ Record.AddSourceLocation(C->getLParenLoc());
+ for (auto *E : C->varlists())
+ Record.AddStmt(E);
+ for (auto *D : C->all_decls())
+ Record.AddDeclRef(D);
+ for (auto N : C->all_num_lists())
+ Record.push_back(N);
+ for (auto N : C->all_lists_sizes())
+ Record.push_back(N);
+ for (auto &M : C->all_components()) {
+ Record.AddStmt(M.getAssociatedExpression());
+ Record.AddDeclRef(M.getAssociatedDeclaration());
+ }
+}
+
void OMPClauseWriter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
Record.push_back(C->varlist_size());
Record.push_back(C->getUniqueDeclarationsNum());
Index: clang/lib/Serialization/ASTReader.cpp
===================================================================
--- clang/lib/Serialization/ASTReader.cpp
+++ clang/lib/Serialization/ASTReader.cpp
@@ -11918,6 +11918,15 @@
C = OMPUseDevicePtrClause::CreateEmpty(Context, Sizes);
break;
}
+ case llvm::omp::OMPC_use_device_addr: {
+ OMPMappableExprListSizeTy Sizes;
+ Sizes.NumVars = Record.readInt();
+ Sizes.NumUniqueDeclarations = Record.readInt();
+ Sizes.NumComponentLists = Record.readInt();
+ Sizes.NumComponents = Record.readInt();
+ C = OMPUseDeviceAddrClause::CreateEmpty(Context, Sizes);
+ break;
+ }
case llvm::omp::OMPC_is_device_ptr: {
OMPMappableExprListSizeTy Sizes;
Sizes.NumVars = Record.readInt();
@@ -12704,6 +12713,48 @@
C->setComponents(Components, ListSizes);
}
+void OMPClauseReader::VisitOMPUseDeviceAddrClause(OMPUseDeviceAddrClause *C) {
+ C->setLParenLoc(Record.readSourceLocation());
+ auto NumVars = C->varlist_size();
+ auto UniqueDecls = C->getUniqueDeclarationsNum();
+ auto TotalLists = C->getTotalComponentListNum();
+ auto TotalComponents = C->getTotalComponentsNum();
+
+ SmallVector<Expr *, 16> Vars;
+ Vars.reserve(NumVars);
+ for (unsigned i = 0; i != NumVars; ++i)
+ Vars.push_back(Record.readSubExpr());
+ C->setVarRefs(Vars);
+
+ SmallVector<ValueDecl *, 16> Decls;
+ Decls.reserve(UniqueDecls);
+ for (unsigned i = 0; i < UniqueDecls; ++i)
+ Decls.push_back(Record.readDeclAs<ValueDecl>());
+ C->setUniqueDecls(Decls);
+
+ SmallVector<unsigned, 16> ListsPerDecl;
+ ListsPerDecl.reserve(UniqueDecls);
+ for (unsigned i = 0; i < UniqueDecls; ++i)
+ ListsPerDecl.push_back(Record.readInt());
+ C->setDeclNumLists(ListsPerDecl);
+
+ SmallVector<unsigned, 32> ListSizes;
+ ListSizes.reserve(TotalLists);
+ for (unsigned i = 0; i < TotalLists; ++i)
+ ListSizes.push_back(Record.readInt());
+ C->setComponentListSizes(ListSizes);
+
+ SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
+ Components.reserve(TotalComponents);
+ for (unsigned i = 0; i < TotalComponents; ++i) {
+ Expr *AssociatedExpr = Record.readSubExpr();
+ auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
+ Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
+ AssociatedExpr, AssociatedDecl));
+ }
+ C->setComponents(Components, ListSizes);
+}
+
void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
C->setLParenLoc(Record.readSourceLocation());
auto NumVars = C->varlist_size();
Index: clang/lib/Sema/TreeTransform.h
===================================================================
--- clang/lib/Sema/TreeTransform.h
+++ clang/lib/Sema/TreeTransform.h
@@ -2036,6 +2036,15 @@
return getSema().ActOnOpenMPUseDevicePtrClause(VarList, Locs);
}
+ /// Build a new OpenMP 'use_device_addr' clause.
+ ///
+ /// By default, performs semantic analysis to build the new OpenMP clause.
+ /// Subclasses may override this routine to provide different behavior.
+ OMPClause *RebuildOMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
+ const OMPVarListLocTy &Locs) {
+ return getSema().ActOnOpenMPUseDeviceAddrClause(VarList, Locs);
+ }
+
/// Build a new OpenMP 'is_device_ptr' clause.
///
/// By default, performs semantic analysis to build the new OpenMP clause.
@@ -9741,6 +9750,21 @@
}
template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPUseDeviceAddrClause(
+ OMPUseDeviceAddrClause *C) {
+ llvm::SmallVector<Expr *, 16> Vars;
+ Vars.reserve(C->varlist_size());
+ for (auto *VE : C->varlists()) {
+ ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+ if (EVar.isInvalid())
+ return nullptr;
+ Vars.push_back(EVar.get());
+ }
+ OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
+ return getDerived().RebuildOMPUseDeviceAddrClause(Vars, Locs);
+}
+
+template <typename Derived>
OMPClause *
TreeTransform<Derived>::TransformOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
llvm::SmallVector<Expr *, 16> Vars;
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -5408,6 +5408,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_nontemporal:
case OMPC_order:
@@ -10167,10 +10168,15 @@
// OpenMP [2.10.1, Restrictions, p. 97]
// At least one map clause must appear on the directive.
- if (!hasClauses(Clauses, OMPC_map, OMPC_use_device_ptr)) {
+ if (!hasClauses(Clauses, OMPC_map, OMPC_use_device_ptr) &&
+ (LangOpts.OpenMP < 50 || !hasClauses(Clauses, OMPC_use_device_addr))) {
+ StringRef Expected;
+ if (LangOpts.OpenMP < 50)
+ Expected = "'map' or 'use_device_ptr'";
+ else
+ Expected = "'map', 'use_device_ptr', or 'use_device_addr'";
Diag(StartLoc, diag::err_omp_no_clause_for_directive)
- << "'map' or 'use_device_ptr'"
- << getOpenMPDirectiveName(OMPD_target_data);
+ << Expected << getOpenMPDirectiveName(OMPD_target_data);
return StmtError();
}
@@ -11535,6 +11541,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -12289,6 +12296,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -12731,6 +12739,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -12956,6 +12965,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -13195,6 +13205,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_atomic_default_mem_order:
case OMPC_device_type:
@@ -13406,6 +13417,9 @@
case OMPC_use_device_ptr:
Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs);
break;
+ case OMPC_use_device_addr:
+ Res = ActOnOpenMPUseDeviceAddrClause(VarList, Locs);
+ break;
case OMPC_is_device_ptr:
Res = ActOnOpenMPIsDevicePtrClause(VarList, Locs);
break;
@@ -18389,6 +18403,54 @@
MVLI.VarBaseDeclarations, MVLI.VarComponents);
}
+OMPClause *Sema::ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
+ const OMPVarListLocTy &Locs) {
+ MappableVarListInfo MVLI(VarList);
+
+ for (Expr *RefExpr : VarList) {
+ assert(RefExpr && "NULL expr in OpenMP use_device_addr clause.");
+ SourceLocation ELoc;
+ SourceRange ERange;
+ Expr *SimpleRefExpr = RefExpr;
+ auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange,
+ /*AllowArraySection=*/true);
+ if (Res.second) {
+ // It will be analyzed later.
+ MVLI.ProcessedVarList.push_back(RefExpr);
+ }
+ ValueDecl *D = Res.first;
+ if (!D)
+ continue;
+ auto *VD = dyn_cast<VarDecl>(D);
+
+ // If required, build a capture to implement the privatization initialized
+ // with the current list item value.
+ DeclRefExpr *Ref = nullptr;
+ if (!VD)
+ Ref = buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+ MVLI.ProcessedVarList.push_back(VD ? RefExpr->IgnoreParens() : Ref);
+
+ // We need to add a data sharing attribute for this variable to make sure it
+ // is correctly captured. A variable that shows up in a use_device_addr has
+ // similar properties of a first private variable.
+ DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);
+
+ // Create a mappable component for the list item. List items in this clause
+ // only need a component.
+ MVLI.VarBaseDeclarations.push_back(D);
+ MVLI.VarComponents.emplace_back();
+ MVLI.VarComponents.back().push_back(
+ OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+ }
+
+ if (MVLI.ProcessedVarList.empty())
+ return nullptr;
+
+ return OMPUseDeviceAddrClause::Create(Context, Locs, MVLI.ProcessedVarList,
+ MVLI.VarBaseDeclarations,
+ MVLI.VarComponents);
+}
+
OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs) {
MappableVarListInfo MVLI(VarList);
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -2497,7 +2497,7 @@
/// in_reduction-clause | allocator-clause | allocate-clause |
/// acq_rel-clause | acquire-clause | release-clause | relaxed-clause |
/// depobj-clause | destroy-clause | detach-clause | inclusive-clause |
-/// exclusive-clause | uses_allocators-clause
+/// exclusive-clause | uses_allocators-clause | use_device_addr-clause
///
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
OpenMPClauseKind CKind, bool FirstClause) {
@@ -2663,6 +2663,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_allocate:
case OMPC_nontemporal:
@@ -3581,6 +3582,8 @@
/// 'from' '(' [ mapper '(' mapper-identifier ')' ':' ] list ')'
/// use_device_ptr-clause:
/// 'use_device_ptr' '(' list ')'
+/// use_device_addr-clause:
+/// 'use_device_addr' '(' list ')'
/// is_device_ptr-clause:
/// 'is_device_ptr' '(' list ')'
/// allocate-clause:
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -4730,6 +4730,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
Index: clang/lib/Basic/OpenMPKinds.cpp
===================================================================
--- clang/lib/Basic/OpenMPKinds.cpp
+++ clang/lib/Basic/OpenMPKinds.cpp
@@ -163,6 +163,7 @@
case OMPC_hint:
case OMPC_uniform:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -411,6 +412,7 @@
case OMPC_hint:
case OMPC_uniform:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
Index: clang/lib/AST/StmtProfile.cpp
===================================================================
--- clang/lib/AST/StmtProfile.cpp
+++ clang/lib/AST/StmtProfile.cpp
@@ -784,6 +784,10 @@
const OMPUseDevicePtrClause *C) {
VisitOMPClauseList(C);
}
+void OMPClauseProfiler::VisitOMPUseDeviceAddrClause(
+ const OMPUseDeviceAddrClause *C) {
+ VisitOMPClauseList(C);
+}
void OMPClauseProfiler::VisitOMPIsDevicePtrClause(
const OMPIsDevicePtrClause *C) {
VisitOMPClauseList(C);
Index: clang/lib/AST/OpenMPClause.cpp
===================================================================
--- clang/lib/AST/OpenMPClause.cpp
+++ clang/lib/AST/OpenMPClause.cpp
@@ -136,6 +136,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -227,6 +228,7 @@
case OMPC_to:
case OMPC_from:
case OMPC_use_device_ptr:
+ case OMPC_use_device_addr:
case OMPC_is_device_ptr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
@@ -1198,6 +1200,53 @@
return new (Mem) OMPUseDevicePtrClause(Sizes);
}
+OMPUseDeviceAddrClause *
+OMPUseDeviceAddrClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> Vars,
+ ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists) {
+ OMPMappableExprListSizeTy Sizes;
+ Sizes.NumVars = Vars.size();
+ Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
+ Sizes.NumComponentLists = ComponentLists.size();
+ Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
+
+ // We need to allocate:
+ // 3 x NumVars x Expr* - we have an original list expression for each clause
+ // list entry and an equal number of private copies and inits.
+ // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
+ // with each component list.
+ // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
+ // number of lists for each unique declaration and the size of each component
+ // list.
+ // NumComponents x MappableComponent - the total of all the components in all
+ // the lists.
+ void *Mem = C.Allocate(
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent>(
+ Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+ Sizes.NumComponents));
+
+ auto *Clause = new (Mem) OMPUseDeviceAddrClause(Locs, Sizes);
+
+ Clause->setVarRefs(Vars);
+ Clause->setClauseInfo(Declarations, ComponentLists);
+ return Clause;
+}
+
+OMPUseDeviceAddrClause *
+OMPUseDeviceAddrClause::CreateEmpty(const ASTContext &C,
+ const OMPMappableExprListSizeTy &Sizes) {
+ void *Mem = C.Allocate(
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent>(
+ Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+ Sizes.NumComponents));
+ return new (Mem) OMPUseDeviceAddrClause(Sizes);
+}
+
OMPIsDevicePtrClause *
OMPIsDevicePtrClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> Vars,
@@ -1934,6 +1983,15 @@
}
}
+void OMPClausePrinter::VisitOMPUseDeviceAddrClause(
+ OMPUseDeviceAddrClause *Node) {
+ if (!Node->varlist_empty()) {
+ OS << "use_device_addr";
+ VisitOMPClauseList(Node, '(');
+ OS << ")";
+ }
+}
+
void OMPClausePrinter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *Node) {
if (!Node->varlist_empty()) {
OS << "is_device_ptr";
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -10776,6 +10776,9 @@
/// Called on well-formed 'use_device_ptr' clause.
OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs);
+ /// Called on well-formed 'use_device_addr' clause.
+ OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
+ const OMPVarListLocTy &Locs);
/// Called on well-formed 'is_device_ptr' clause.
OMPClause *ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs);
Index: clang/include/clang/AST/RecursiveASTVisitor.h
===================================================================
--- clang/include/clang/AST/RecursiveASTVisitor.h
+++ clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3522,6 +3522,13 @@
}
template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPUseDeviceAddrClause(
+ OMPUseDeviceAddrClause *C) {
+ TRY_TO(VisitOMPClauseList(C));
+ return true;
+}
+
+template <typename Derived>
bool RecursiveASTVisitor<Derived>::VisitOMPIsDevicePtrClause(
OMPIsDevicePtrClause *C) {
TRY_TO(VisitOMPClauseList(C));
Index: clang/include/clang/AST/OpenMPClause.h
===================================================================
--- clang/include/clang/AST/OpenMPClause.h
+++ clang/include/clang/AST/OpenMPClause.h
@@ -6597,6 +6597,110 @@
}
};
+/// This represents clause 'use_device_addr' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target data use_device_addr(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target data' has clause
+/// 'use_device_addr' with the variables 'a' and 'b'.
+class OMPUseDeviceAddrClause final
+ : public OMPMappableExprListClause<OMPUseDeviceAddrClause>,
+ private llvm::TrailingObjects<
+ OMPUseDeviceAddrClause, Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent> {
+ friend class OMPClauseReader;
+ friend OMPMappableExprListClause;
+ friend OMPVarListClause;
+ friend TrailingObjects;
+
+ /// Build clause with number of variables \a NumVars.
+ ///
+ /// \param Locs Locations needed to build a mappable clause. It includes 1)
+ /// StartLoc: starting location of the clause (the clause keyword); 2)
+ /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
+ /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+ /// NumVars: number of expressions listed in this clause; 2)
+ /// NumUniqueDeclarations: number of unique base declarations in this clause;
+ /// 3) NumComponentLists: number of component lists in this clause; and 4)
+ /// NumComponents: total number of expression components in the clause.
+ explicit OMPUseDeviceAddrClause(const OMPVarListLocTy &Locs,
+ const OMPMappableExprListSizeTy &Sizes)
+ : OMPMappableExprListClause(llvm::omp::OMPC_use_device_addr, Locs,
+ Sizes) {}
+
+ /// Build an empty clause.
+ ///
+ /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+ /// NumVars: number of expressions listed in this clause; 2)
+ /// NumUniqueDeclarations: number of unique base declarations in this clause;
+ /// 3) NumComponentLists: number of component lists in this clause; and 4)
+ /// NumComponents: total number of expression components in the clause.
+ explicit OMPUseDeviceAddrClause(const OMPMappableExprListSizeTy &Sizes)
+ : OMPMappableExprListClause(llvm::omp::OMPC_use_device_addr,
+ OMPVarListLocTy(), Sizes) {}
+
+ /// Define the sizes of each trailing object array except the last one. This
+ /// is required for TrailingObjects to work properly.
+ size_t numTrailingObjects(OverloadToken<Expr *>) const {
+ return varlist_size();
+ }
+ size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
+ return getUniqueDeclarationsNum();
+ }
+ size_t numTrailingObjects(OverloadToken<unsigned>) const {
+ return getUniqueDeclarationsNum() + getTotalComponentListNum();
+ }
+
+public:
+ /// Creates clause with a list of variables \a Vars.
+ ///
+ /// \param C AST context.
+ /// \param Locs Locations needed to build a mappable clause. It includes 1)
+ /// StartLoc: starting location of the clause (the clause keyword); 2)
+ /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
+ /// \param Vars The original expression used in the clause.
+ /// \param Declarations Declarations used in the clause.
+ /// \param ComponentLists Component lists used in the clause.
+ static OMPUseDeviceAddrClause *
+ Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists);
+
+ /// Creates an empty clause with the place for \a NumVars variables.
+ ///
+ /// \param C AST context.
+ /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+ /// NumVars: number of expressions listed in this clause; 2)
+ /// NumUniqueDeclarations: number of unique base declarations in this clause;
+ /// 3) NumComponentLists: number of component lists in this clause; and 4)
+ /// NumComponents: total number of expression components in the clause.
+ static OMPUseDeviceAddrClause *
+ CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes);
+
+ child_range children() {
+ return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+ reinterpret_cast<Stmt **>(varlist_end()));
+ }
+
+ const_child_range children() const {
+ auto Children = const_cast<OMPUseDeviceAddrClause *>(this)->children();
+ return const_child_range(Children.begin(), Children.end());
+ }
+
+ child_range used_children() {
+ return child_range(child_iterator(), child_iterator());
+ }
+ const_child_range used_children() const {
+ return const_child_range(const_child_iterator(), const_child_iterator());
+ }
+
+ static bool classof(const OMPClause *T) {
+ return T->getClauseKind() == llvm::omp::OMPC_use_device_addr;
+ }
+};
+
/// This represents clause 'is_device_ptr' in the '#pragma omp ...'
/// directives.
///
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
