<div dir="ltr">Samuel,<div><br></div><div>I think your change breaks the build:</div><div><br></div><div><a href="http://lab.llvm.org:8011/builders/sanitizer-windows/builds/26397">http://lab.llvm.org:8011/builders/sanitizer-windows/builds/26397</a><br></div><div><br></div><div><div>FAILED: tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/CGOpenMPRuntime.cpp.obj </div><div>C:\PROGRA~2\MICROS~1.0\VC\bin\AMD64_~1\cl.exe /nologo /TP -DCLANG_ENABLE_ARCMT -DCLANG_ENABLE_OBJC_REWRITER -DCLANG_ENABLE_STATIC_ANALYZER -DGTEST_HAS_RTTI=0 -DUNICODE -D_CRT_NONSTDC_NO_DEPRECATE -D_CRT_NONSTDC_NO_WARNINGS -D_CRT_SECURE_NO_DEPRECATE -D_CRT_SECURE_NO_WARNINGS -D_DEBUG_POINTER_IMPL="" -D_GNU_SOURCE -D_HAS_EXCEPTIONS=0 -D_SCL_SECURE_NO_DEPRECATE -D_SCL_SECURE_NO_WARNINGS -D_UNICODE -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -Itools\clang\lib\CodeGen -IC:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen -IC:\b\slave\sanitizer-windows\llvm\tools\clang\include -Itools\clang\include -Iinclude -IC:\b\slave\sanitizer-windows\llvm\include /DWIN32 /D_WINDOWS /W4 -wd4141 -wd4146 -wd4180 -wd4244 -wd4258 -wd4267 -wd4291 -wd4345 -wd4351 -wd4355 -wd4456 -wd4457 -wd4458 -wd4459 -wd4503 -wd4624 -wd4722 -wd4800 -wd4100 -wd4127 -wd4512 -wd4505 -wd4610 -wd4510 -wd4702 -wd4245 -wd4706 -wd4310 -wd4701 -wd4703 -wd4389 -wd4611 -wd4805 -wd4204 -wd4577 -wd4091 -wd4592 -wd4319 -wd4324 -w14062 -we4238 /Zc:inline /Oi /Zc:rvalueCast /MD /O2 /Ob2 -UNDEBUG /EHs-c- /GR- /showIncludes /Fotools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\CGOpenMPRuntime.cpp.obj /Fdtools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\ /FS -c C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5482) : error C3486: a parameter for a lambda cannot have a default argument</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is not a type name, static, or enumerator</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) : error C2065: 'Directive' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) : error C2228: left of '.getClausesOfKind' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) : error C2059: syntax error : ')'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) : error C2143: syntax error : missing ';' before '<end Parse>'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is not a type name, static, or enumerator</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) : error C2065: 'Directive' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) : error C2228: left of '.getClausesOfKind' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) : error C2059: syntax error : ')'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) : error C2143: syntax error : missing ';' before '<end Parse>'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5493) : error C2143: syntax error : missing ')' before ';'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is not a type name, static, or enumerator</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) : error C2065: 'Directive' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) : error C2228: left of '.getClausesOfKind' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) : error C2059: syntax error : ')'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) : error C2143: syntax error : missing ';' before '<end Parse>'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5496) : error C2143: syntax error : missing ')' before ';'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is not a type name, static, or enumerator</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) : error C2065: 'Directive' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) : error C2228: left of '.getClausesOfKind' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) : error C2059: syntax error : ')'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) : error C2143: syntax error : missing ';' before '<end Parse>'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5503) : error C2143: syntax error : missing ';' before '{'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5503) : error C2143: syntax error : missing ')' before ';'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504) : error C2065: 'L' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504) : error C2228: left of '.second' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504) : error C2228: left of '.empty' must have class/struct/union</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) : error C2065: 'L' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) : error C2228: left of '.second' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) : error C2228: left of '.back' must have class/struct/union</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) : error C2228: left of '.getAssociatedDeclaration' must have class/struct/union</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) : error C2065: 'L' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) : error C2228: left of '.second' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) : error C2228: left of '.back' must have class/struct/union</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) : error C2228: left of '.getAssociatedExpression' must have class/struct/union</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5511) : error C3536: 'IE': cannot be used before it is initialized</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5523) : error C3536: 'IE': cannot be used before it is initialized</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5528) : error C2044: illegal continue</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) : error C2327: '`anonymous-namespace'::MappableExprsHandler::CGF' : is not a type name, static, or enumerator</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) : error C2065: 'CGF' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) : error C2228: left of '.EmitLoadOfLValue' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) : error C2228: left of '.EmitLValue' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) : error C3536: 'IE': cannot be used before it is initialized</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5536) : error C2228: left of '.getScalarVal' must have class/struct/union</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539) : error C2327: '`anonymous-namespace'::MappableExprsHandler::CGF' : is not a type name, static, or enumerator</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539) : error C2065: 'CGF' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539) : error C2228: left of '.SizeTy' must have class/struct/union</div><div> type is 'unknown-type'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) : error C2143: syntax error : missing ')' before ';'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) : error C2143: syntax error : missing ';' before ')'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) : error C2065: '$S239' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) : error C2065: '$S240' : undeclared identifier</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) : error C2143: syntax error : missing ';' before '{'</div><div>C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) : fatal error C1903: unable to recover from previous error(s); stopping compilation</div><div>18446743528493.273 [136/2/116] Building CXX object tools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\CGOpenCLRuntime.cpp.obj</div></div></div><br><div class="gmail_quote"><div dir="ltr">On Thu, Jul 28, 2016 at 7:31 AM Samuel Antao via cfe-commits <<a href="mailto:cfe-commits@lists.llvm.org">cfe-commits@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Author: sfantao<br>
Date: Thu Jul 28 09:23:26 2016<br>
New Revision: 276977<br>
<br>
URL: <a href="http://llvm.org/viewvc/llvm-project?rev=276977&view=rev" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project?rev=276977&view=rev</a><br>
Log:<br>
[OpenMP] Codegen for use_device_ptr clause.<br>
<br>
Summary: This patch adds support for the use_device_ptr clause. It includes changes in SEMA that could not be tested without codegen, namely, the use of the first private logic and mappable expressions support.<br>
<br>
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev<br>
<br>
Subscribers: caomhin, cfe-commits<br>
<br>
Differential Revision: <a href="https://reviews.llvm.org/D22691" rel="noreferrer" target="_blank">https://reviews.llvm.org/D22691</a><br>
<br>
Added:<br>
cfe/trunk/test/OpenMP/target_data_use_device_ptr_codegen.cpp<br>
Modified:<br>
cfe/trunk/include/clang/AST/OpenMPClause.h<br>
cfe/trunk/lib/AST/OpenMPClause.cpp<br>
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp<br>
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h<br>
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp<br>
cfe/trunk/lib/CodeGen/CodeGenFunction.h<br>
cfe/trunk/lib/Sema/SemaOpenMP.cpp<br>
cfe/trunk/lib/Serialization/ASTReaderStmt.cpp<br>
cfe/trunk/lib/Serialization/ASTWriterStmt.cpp<br>
<br>
Modified: cfe/trunk/include/clang/AST/OpenMPClause.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/OpenMPClause.h?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/OpenMPClause.h?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/include/clang/AST/OpenMPClause.h (original)<br>
+++ cfe/trunk/include/clang/AST/OpenMPClause.h Thu Jul 28 09:23:26 2016<br>
@@ -4228,50 +4228,153 @@ public:<br>
/// 'use_device_ptr' with the variables 'a' and 'b'.<br>
///<br>
class OMPUseDevicePtrClause final<br>
- : public OMPVarListClause<OMPUseDevicePtrClause>,<br>
- private llvm::TrailingObjects<OMPUseDevicePtrClause, Expr *> {<br>
+ : public OMPMappableExprListClause<OMPUseDevicePtrClause>,<br>
+ private llvm::TrailingObjects<<br>
+ OMPUseDevicePtrClause, Expr *, ValueDecl *, unsigned,<br>
+ OMPClauseMappableExprCommon::MappableComponent> {<br>
friend TrailingObjects;<br>
friend OMPVarListClause;<br>
+ friend OMPMappableExprListClause;<br>
friend class OMPClauseReader;<br>
- /// Build clause with number of variables \a N.<br>
+<br>
+ /// Define the sizes of each trailing object array except the last one. This<br>
+ /// is required for TrailingObjects to work properly.<br>
+ size_t numTrailingObjects(OverloadToken<Expr *>) const {<br>
+ return 3 * varlist_size();<br>
+ }<br>
+ size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {<br>
+ return getUniqueDeclarationsNum();<br>
+ }<br>
+ size_t numTrailingObjects(OverloadToken<unsigned>) const {<br>
+ return getUniqueDeclarationsNum() + getTotalComponentListNum();<br>
+ }<br>
+<br>
+ /// Build clause with number of variables \a NumVars.<br>
///<br>
/// \param StartLoc Starting location of the clause.<br>
- /// \param LParenLoc Location of '('.<br>
/// \param EndLoc Ending location of the clause.<br>
- /// \param N Number of the variables in the clause.<br>
- ///<br>
- OMPUseDevicePtrClause(SourceLocation StartLoc, SourceLocation LParenLoc,<br>
- SourceLocation EndLoc, unsigned N)<br>
- : OMPVarListClause<OMPUseDevicePtrClause>(OMPC_use_device_ptr, StartLoc,<br>
- LParenLoc, EndLoc, N) {}<br>
+ /// \param NumVars Number of expressions listed in this clause.<br>
+ /// \param NumUniqueDeclarations Number of unique base declarations in this<br>
+ /// clause.<br>
+ /// \param NumComponentLists Number of component lists in this clause.<br>
+ /// \param NumComponents Total number of expression components in the clause.<br>
+ ///<br>
+ explicit OMPUseDevicePtrClause(SourceLocation StartLoc,<br>
+ SourceLocation LParenLoc,<br>
+ SourceLocation EndLoc, unsigned NumVars,<br>
+ unsigned NumUniqueDeclarations,<br>
+ unsigned NumComponentLists,<br>
+ unsigned NumComponents)<br>
+ : OMPMappableExprListClause(OMPC_use_device_ptr, StartLoc, LParenLoc,<br>
+ EndLoc, NumVars, NumUniqueDeclarations,<br>
+ NumComponentLists, NumComponents) {}<br>
<br>
- /// \brief Build an empty clause.<br>
- ///<br>
- /// \param N Number of variables.<br>
+ /// Build an empty clause.<br>
///<br>
- explicit OMPUseDevicePtrClause(unsigned N)<br>
- : OMPVarListClause<OMPUseDevicePtrClause>(<br>
- OMPC_use_device_ptr, SourceLocation(), SourceLocation(),<br>
- SourceLocation(), N) {}<br>
+ /// \param NumVars Number of expressions listed in this clause.<br>
+ /// \param NumUniqueDeclarations Number of unique base declarations in this<br>
+ /// clause.<br>
+ /// \param NumComponentLists Number of component lists in this clause.<br>
+ /// \param NumComponents Total number of expression components in the clause.<br>
+ ///<br>
+ explicit OMPUseDevicePtrClause(unsigned NumVars,<br>
+ unsigned NumUniqueDeclarations,<br>
+ unsigned NumComponentLists,<br>
+ unsigned NumComponents)<br>
+ : OMPMappableExprListClause(OMPC_use_device_ptr, SourceLocation(),<br>
+ SourceLocation(), SourceLocation(), NumVars,<br>
+ NumUniqueDeclarations, NumComponentLists,<br>
+ NumComponents) {}<br>
+<br>
+ /// Sets the list of references to private copies with initializers for new<br>
+ /// private variables.<br>
+ /// \param VL List of references.<br>
+ void setPrivateCopies(ArrayRef<Expr *> VL);<br>
+<br>
+ /// Gets the list of references to private copies with initializers for new<br>
+ /// private variables.<br>
+ MutableArrayRef<Expr *> getPrivateCopies() {<br>
+ return MutableArrayRef<Expr *>(varlist_end(), varlist_size());<br>
+ }<br>
+ ArrayRef<const Expr *> getPrivateCopies() const {<br>
+ return llvm::makeArrayRef(varlist_end(), varlist_size());<br>
+ }<br>
+<br>
+ /// Sets the list of references to initializer variables for new private<br>
+ /// variables.<br>
+ /// \param VL List of references.<br>
+ void setInits(ArrayRef<Expr *> VL);<br>
+<br>
+ /// Gets the list of references to initializer variables for new private<br>
+ /// variables.<br>
+ MutableArrayRef<Expr *> getInits() {<br>
+ return MutableArrayRef<Expr *>(getPrivateCopies().end(), varlist_size());<br>
+ }<br>
+ ArrayRef<const Expr *> getInits() const {<br>
+ return llvm::makeArrayRef(getPrivateCopies().end(), varlist_size());<br>
+ }<br>
<br>
public:<br>
- /// Creates clause with a list of variables \a VL.<br>
+ /// Creates clause with a list of variables \a Vars.<br>
///<br>
/// \param C AST context.<br>
/// \param StartLoc Starting location of the clause.<br>
- /// \param LParenLoc Location of '('.<br>
/// \param EndLoc Ending location of the clause.<br>
- /// \param VL List of references to the variables.<br>
+ /// \param Vars The original expression used in the clause.<br>
+ /// \param PrivateVars Expressions referring to private copies.<br>
+ /// \param Inits Expressions referring to private copy initializers.<br>
+ /// \param Declarations Declarations used in the clause.<br>
+ /// \param ComponentLists Component lists used in the clause.<br>
///<br>
static OMPUseDevicePtrClause *<br>
Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,<br>
- SourceLocation EndLoc, ArrayRef<Expr *> VL);<br>
- /// Creates an empty clause with the place for \a N variables.<br>
+ SourceLocation EndLoc, ArrayRef<Expr *> Vars,<br>
+ ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits,<br>
+ ArrayRef<ValueDecl *> Declarations,<br>
+ MappableExprComponentListsRef ComponentLists);<br>
+<br>
+ /// Creates an empty clause with the place for \a NumVars variables.<br>
///<br>
/// \param C AST context.<br>
- /// \param N The number of variables.<br>
- ///<br>
- static OMPUseDevicePtrClause *CreateEmpty(const ASTContext &C, unsigned N);<br>
+ /// \param NumVars Number of expressions listed in the clause.<br>
+ /// \param NumUniqueDeclarations Number of unique base declarations in this<br>
+ /// clause.<br>
+ /// \param NumComponentLists Number of unique base declarations in this<br>
+ /// clause.<br>
+ /// \param NumComponents Total number of expression components in the clause.<br>
+ ///<br>
+ static OMPUseDevicePtrClause *CreateEmpty(const ASTContext &C,<br>
+ unsigned NumVars,<br>
+ unsigned NumUniqueDeclarations,<br>
+ unsigned NumComponentLists,<br>
+ unsigned NumComponents);<br>
+<br>
+ typedef MutableArrayRef<Expr *>::iterator private_copies_iterator;<br>
+ typedef ArrayRef<const Expr *>::iterator private_copies_const_iterator;<br>
+ typedef llvm::iterator_range<private_copies_iterator> private_copies_range;<br>
+ typedef llvm::iterator_range<private_copies_const_iterator><br>
+ private_copies_const_range;<br>
+<br>
+ private_copies_range private_copies() {<br>
+ return private_copies_range(getPrivateCopies().begin(),<br>
+ getPrivateCopies().end());<br>
+ }<br>
+ private_copies_const_range private_copies() const {<br>
+ return private_copies_const_range(getPrivateCopies().begin(),<br>
+ getPrivateCopies().end());<br>
+ }<br>
+<br>
+ typedef MutableArrayRef<Expr *>::iterator inits_iterator;<br>
+ typedef ArrayRef<const Expr *>::iterator inits_const_iterator;<br>
+ typedef llvm::iterator_range<inits_iterator> inits_range;<br>
+ typedef llvm::iterator_range<inits_const_iterator> inits_const_range;<br>
+<br>
+ inits_range inits() {<br>
+ return inits_range(getInits().begin(), getInits().end());<br>
+ }<br>
+ inits_const_range inits() const {<br>
+ return inits_const_range(getInits().begin(), getInits().end());<br>
+ }<br>
<br>
child_range children() {<br>
return child_range(reinterpret_cast<Stmt **>(varlist_begin()),<br>
<br>
Modified: cfe/trunk/lib/AST/OpenMPClause.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/AST/OpenMPClause.cpp (original)<br>
+++ cfe/trunk/lib/AST/OpenMPClause.cpp Thu Jul 28 09:23:26 2016<br>
@@ -732,22 +732,66 @@ OMPFromClause *OMPFromClause::CreateEmpt<br>
NumComponentLists, NumComponents);<br>
}<br>
<br>
-OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(const ASTContext &C,<br>
- SourceLocation StartLoc,<br>
- SourceLocation LParenLoc,<br>
- SourceLocation EndLoc,<br>
- ArrayRef<Expr *> VL) {<br>
- void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));<br>
- OMPUseDevicePtrClause *Clause =<br>
- new (Mem) OMPUseDevicePtrClause(StartLoc, LParenLoc, EndLoc, VL.size());<br>
- Clause->setVarRefs(VL);<br>
+void OMPUseDevicePtrClause::setPrivateCopies(ArrayRef<Expr *> VL) {<br>
+ assert(VL.size() == varlist_size() &&<br>
+ "Number of private copies is not the same as the preallocated buffer");<br>
+ std::copy(VL.begin(), VL.end(), varlist_end());<br>
+}<br>
+<br>
+void OMPUseDevicePtrClause::setInits(ArrayRef<Expr *> VL) {<br>
+ assert(VL.size() == varlist_size() &&<br>
+ "Number of inits is not the same as the preallocated buffer");<br>
+ std::copy(VL.begin(), VL.end(), getPrivateCopies().end());<br>
+}<br>
+<br>
+OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(<br>
+ const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,<br>
+ SourceLocation EndLoc, ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars,<br>
+ ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations,<br>
+ MappableExprComponentListsRef ComponentLists) {<br>
+ unsigned NumVars = Vars.size();<br>
+ unsigned NumUniqueDeclarations =<br>
+ getUniqueDeclarationsTotalNumber(Declarations);<br>
+ unsigned NumComponentLists = ComponentLists.size();<br>
+ unsigned NumComponents = getComponentsTotalNumber(ComponentLists);<br>
+<br>
+ // We need to allocate:<br>
+ // 3 x NumVars x Expr* - we have an original list expression for each clause<br>
+ // list entry and an equal number of private copies and inits.<br>
+ // NumUniqueDeclarations x ValueDecl* - unique base declarations associated<br>
+ // with each component list.<br>
+ // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the<br>
+ // number of lists for each unique declaration and the size of each component<br>
+ // list.<br>
+ // NumComponents x MappableComponent - the total of all the components in all<br>
+ // the lists.<br>
+ void *Mem = C.Allocate(<br>
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,<br>
+ OMPClauseMappableExprCommon::MappableComponent>(<br>
+ 3 * NumVars, NumUniqueDeclarations,<br>
+ NumUniqueDeclarations + NumComponentLists, NumComponents));<br>
+<br>
+ OMPUseDevicePtrClause *Clause = new (Mem) OMPUseDevicePtrClause(<br>
+ StartLoc, LParenLoc, EndLoc, NumVars, NumUniqueDeclarations,<br>
+ NumComponentLists, NumComponents);<br>
+<br>
+ Clause->setVarRefs(Vars);<br>
+ Clause->setPrivateCopies(PrivateVars);<br>
+ Clause->setInits(Inits);<br>
+ Clause->setClauseInfo(Declarations, ComponentLists);<br>
return Clause;<br>
}<br>
<br>
-OMPUseDevicePtrClause *OMPUseDevicePtrClause::CreateEmpty(const ASTContext &C,<br>
- unsigned N) {<br>
- void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));<br>
- return new (Mem) OMPUseDevicePtrClause(N);<br>
+OMPUseDevicePtrClause *OMPUseDevicePtrClause::CreateEmpty(<br>
+ const ASTContext &C, unsigned NumVars, unsigned NumUniqueDeclarations,<br>
+ unsigned NumComponentLists, unsigned NumComponents) {<br>
+ void *Mem = C.Allocate(<br>
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,<br>
+ OMPClauseMappableExprCommon::MappableComponent>(<br>
+ 3 * NumVars, NumUniqueDeclarations,<br>
+ NumUniqueDeclarations + NumComponentLists, NumComponents));<br>
+ return new (Mem) OMPUseDevicePtrClause(NumVars, NumUniqueDeclarations,<br>
+ NumComponentLists, NumComponents);<br>
}<br>
<br>
OMPIsDevicePtrClause *OMPIsDevicePtrClause::Create(const ASTContext &C,<br>
<br>
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)<br>
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Thu Jul 28 09:23:26 2016<br>
@@ -4981,6 +4981,9 @@ public:<br>
/// map/privatization results in multiple arguments passed to the runtime<br>
/// library.<br>
OMP_MAP_FIRST_REF = 0x20,<br>
+ /// \brief Signal that the runtime library has to return the device pointer<br>
+ /// in the current position for the data being mapped.<br>
+ OMP_MAP_RETURN_PTR = 0x40,<br>
/// \brief This flag signals that the reference being passed is a pointer to<br>
/// private data.<br>
OMP_MAP_PRIVATE_PTR = 0x80,<br>
@@ -4988,6 +4991,24 @@ public:<br>
OMP_MAP_PRIVATE_VAL = 0x100,<br>
};<br>
<br>
+ /// Class that associates information with a base pointer to be passed to the<br>
+ /// runtime library.<br>
+ class BasePointerInfo {<br>
+ /// The base pointer.<br>
+ llvm::Value *Ptr = nullptr;<br>
+ /// The base declaration that refers to this device pointer, or null if<br>
+ /// there is none.<br>
+ const ValueDecl *DevPtrDecl = nullptr;<br>
+<br>
+ public:<br>
+ BasePointerInfo(llvm::Value *Ptr, const ValueDecl *DevPtrDecl = nullptr)<br>
+ : Ptr(Ptr), DevPtrDecl(DevPtrDecl) {}<br>
+ llvm::Value *operator*() const { return Ptr; }<br>
+ const ValueDecl *getDevicePtrDecl() const { return DevPtrDecl; }<br>
+ void setDevicePtrDecl(const ValueDecl *D) { DevPtrDecl = D; }<br>
+ };<br>
+<br>
+ typedef SmallVector<BasePointerInfo, 16> MapBaseValuesArrayTy;<br>
typedef SmallVector<llvm::Value *, 16> MapValuesArrayTy;<br>
typedef SmallVector<unsigned, 16> MapFlagsArrayTy;<br>
<br>
@@ -5129,7 +5150,7 @@ private:<br>
void generateInfoForComponentList(<br>
OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier,<br>
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,<br>
- MapValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,<br>
+ MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,<br>
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,<br>
bool IsFirstComponentList) const {<br>
<br>
@@ -5400,8 +5421,10 @@ public:<br>
}<br>
<br>
/// \brief Generate all the base pointers, section pointers, sizes and map<br>
- /// types for the extracted mappable expressions.<br>
- void generateAllInfo(MapValuesArrayTy &BasePointers,<br>
+ /// types for the extracted mappable expressions. Also, for each item that<br>
+ /// relates with a device pointer, a pair of the relevant declaration and<br>
+ /// index where it occurs is appended to the device pointers info array.<br>
+ void generateAllInfo(MapBaseValuesArrayTy &BasePointers,<br>
MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,<br>
MapFlagsArrayTy &Types) const {<br>
BasePointers.clear();<br>
@@ -5410,9 +5433,28 @@ public:<br>
Types.clear();<br>
<br>
struct MapInfo {<br>
+ /// Kind that defines how a device pointer has to be returned.<br>
+ enum ReturnPointerKind {<br>
+ // Don't have to return any pointer.<br>
+ RPK_None,<br>
+ // Pointer is the base of the declaration.<br>
+ RPK_Base,<br>
+ // Pointer is a member of the base declaration - 'this'<br>
+ RPK_Member,<br>
+ // Pointer is a reference and a member of the base declaration - 'this'<br>
+ RPK_MemberReference,<br>
+ };<br>
OMPClauseMappableExprCommon::MappableExprComponentListRef Components;<br>
- OpenMPMapClauseKind MapType;<br>
- OpenMPMapClauseKind MapTypeModifier;<br>
+ OpenMPMapClauseKind MapType = OMPC_MAP_unknown;<br>
+ OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;<br>
+ ReturnPointerKind ReturnDevicePointer = RPK_None;<br>
+ MapInfo(<br>
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components,<br>
+ OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier,<br>
+ ReturnPointerKind ReturnDevicePointer)<br>
+ : Components(Components), MapType(MapType),<br>
+ MapTypeModifier(MapTypeModifier),<br>
+ ReturnDevicePointer(ReturnDevicePointer) {}<br>
};<br>
<br>
// We have to process the component lists that relate with the same<br>
@@ -5422,14 +5464,15 @@ public:<br>
<br>
// Helper function to fill the information map for the different supported<br>
// clauses.<br>
- auto &&InfoGen =<br>
- [&Info](const ValueDecl *D,<br>
- OMPClauseMappableExprCommon::MappableExprComponentListRef L,<br>
- OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapModifier) {<br>
- const ValueDecl *VD =<br>
- D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;<br>
- Info[VD].push_back({L, MapType, MapModifier});<br>
- };<br>
+ auto &&InfoGen = [&Info](<br>
+ const ValueDecl *D,<br>
+ OMPClauseMappableExprCommon::MappableExprComponentListRef L,<br>
+ OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapModifier,<br>
+ MapInfo::ReturnPointerKind ReturnDevicePointer = MapInfo::RPK_None) {<br>
+ const ValueDecl *VD =<br>
+ D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;<br>
+ Info[VD].push_back({L, MapType, MapModifier, ReturnDevicePointer});<br>
+ };<br>
<br>
for (auto *C : Directive.getClausesOfKind<OMPMapClause>())<br>
for (auto L : C->component_lists())<br>
@@ -5441,6 +5484,51 @@ public:<br>
for (auto L : C->component_lists())<br>
InfoGen(L.first, L.second, OMPC_MAP_from, OMPC_MAP_unknown);<br>
<br>
+ // Look at the use_device_ptr clause information and mark the existing map<br>
+ // entries as such. If there is no map information for an entry in the<br>
+ // use_device_ptr list, we create one with map type 'alloc' and zero size<br>
+ // section. It is the user fault if that was not mapped before.<br>
+ for (auto *C : Directive.getClausesOfKind<OMPUseDevicePtrClause>())<br>
+ for (auto L : C->component_lists()) {<br>
+ assert(!L.second.empty() && "Not expecting empty list of components!");<br>
+ const ValueDecl *VD = L.second.back().getAssociatedDeclaration();<br>
+ VD = cast<ValueDecl>(VD->getCanonicalDecl());<br>
+ auto *IE = L.second.back().getAssociatedExpression();<br>
+ // If the first component is a member expression, we have to look into<br>
+ // 'this', which maps to null in the map of map information. Otherwise<br>
+ // look directly for the information.<br>
+ auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);<br>
+<br>
+ // We potentially have map information for this declaration already.<br>
+ // Look for the first set of components that refer to it.<br>
+ if (It != Info.end()) {<br>
+ auto CI = std::find_if(<br>
+ It->second.begin(), It->second.end(), [VD](const MapInfo &MI) {<br>
+ return MI.Components.back().getAssociatedDeclaration() == VD;<br>
+ });<br>
+ // If we found a map entry, signal that the pointer has to be returned<br>
+ // and move on to the next declaration.<br>
+ if (CI != It->second.end()) {<br>
+ CI->ReturnDevicePointer = isa<MemberExpr>(IE)<br>
+ ? (VD->getType()->isReferenceType()<br>
+ ? MapInfo::RPK_MemberReference<br>
+ : MapInfo::RPK_Member)<br>
+ : MapInfo::RPK_Base;<br>
+ continue;<br>
+ }<br>
+ }<br>
+<br>
+ // We didn't find any match in our map information - generate a zero<br>
+ // size array section.<br>
+ llvm::Value *Ptr =<br>
+ CGF.EmitLoadOfLValue(CGF.EmitLValue(IE), SourceLocation())<br>
+ .getScalarVal();<br>
+ BasePointers.push_back({Ptr, VD});<br>
+ Pointers.push_back(Ptr);<br>
+ Sizes.push_back(llvm::Constant::getNullValue(CGF.SizeTy));<br>
+ Types.push_back(OMP_MAP_RETURN_PTR | OMP_MAP_FIRST_REF);<br>
+ }<br>
+<br>
for (auto &M : Info) {<br>
// We need to know when we generate information for the first component<br>
// associated with a capture, because the mapping flags depend on it.<br>
@@ -5448,9 +5536,35 @@ public:<br>
for (MapInfo &L : M.second) {<br>
assert(!L.Components.empty() &&<br>
"Not expecting declaration with no component lists.");<br>
+<br>
+ // Remember the current base pointer index.<br>
+ unsigned CurrentBasePointersIdx = BasePointers.size();<br>
generateInfoForComponentList(L.MapType, L.MapTypeModifier, L.Components,<br>
BasePointers, Pointers, Sizes, Types,<br>
IsFirstComponentList);<br>
+<br>
+ // If this entry relates with a device pointer, set the relevant<br>
+ // declaration and add the 'return pointer' flag.<br>
+ if (IsFirstComponentList &&<br>
+ L.ReturnDevicePointer != MapInfo::RPK_None) {<br>
+ // If the pointer is not the base of the map, we need to skip the<br>
+ // base. If it is a reference in a member field, we also need to skip<br>
+ // the map of the reference.<br>
+ if (L.ReturnDevicePointer != MapInfo::RPK_Base) {<br>
+ ++CurrentBasePointersIdx;<br>
+ if (L.ReturnDevicePointer == MapInfo::RPK_MemberReference)<br>
+ ++CurrentBasePointersIdx;<br>
+ }<br>
+ assert(BasePointers.size() > CurrentBasePointersIdx &&<br>
+ "Unexpected number of mapped base pointers.");<br>
+<br>
+ auto *RelevantVD = L.Components.back().getAssociatedDeclaration();<br>
+ assert(RelevantVD &&<br>
+ "No relevant declaration related with device pointer??");<br>
+<br>
+ BasePointers[CurrentBasePointersIdx].setDevicePtrDecl(RelevantVD);<br>
+ Types[CurrentBasePointersIdx] |= OMP_MAP_RETURN_PTR;<br>
+ }<br>
IsFirstComponentList = false;<br>
}<br>
}<br>
@@ -5459,7 +5573,7 @@ public:<br>
/// \brief Generate the base pointers, section pointers, sizes and map types<br>
/// associated to a given capture.<br>
void generateInfoForCapture(const CapturedStmt::Capture *Cap,<br>
- MapValuesArrayTy &BasePointers,<br>
+ MapBaseValuesArrayTy &BasePointers,<br>
MapValuesArrayTy &Pointers,<br>
MapValuesArrayTy &Sizes,<br>
MapFlagsArrayTy &Types) const {<br>
@@ -5496,12 +5610,12 @@ public:<br>
<br>
/// \brief Generate the default map information for a given capture \a CI,<br>
/// record field declaration \a RI and captured value \a CV.<br>
- void generateDefaultMapInfo(<br>
- const CapturedStmt::Capture &CI, const FieldDecl &RI, llvm::Value *CV,<br>
- MappableExprsHandler::MapValuesArrayTy &CurBasePointers,<br>
- MappableExprsHandler::MapValuesArrayTy &CurPointers,<br>
- MappableExprsHandler::MapValuesArrayTy &CurSizes,<br>
- MappableExprsHandler::MapFlagsArrayTy &CurMapTypes) {<br>
+ void generateDefaultMapInfo(const CapturedStmt::Capture &CI,<br>
+ const FieldDecl &RI, llvm::Value *CV,<br>
+ MapBaseValuesArrayTy &CurBasePointers,<br>
+ MapValuesArrayTy &CurPointers,<br>
+ MapValuesArrayTy &CurSizes,<br>
+ MapFlagsArrayTy &CurMapTypes) {<br>
<br>
// Do the default mapping.<br>
if (CI.capturesThis()) {<br>
@@ -5510,15 +5624,14 @@ public:<br>
const PointerType *PtrTy = cast<PointerType>(RI.getType().getTypePtr());<br>
CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType()));<br>
// Default map type.<br>
- CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO |<br>
- MappableExprsHandler::OMP_MAP_FROM);<br>
+ CurMapTypes.push_back(OMP_MAP_TO | OMP_MAP_FROM);<br>
} else if (CI.capturesVariableByCopy()) {<br>
CurBasePointers.push_back(CV);<br>
CurPointers.push_back(CV);<br>
if (!RI.getType()->isAnyPointerType()) {<br>
// We have to signal to the runtime captures passed by value that are<br>
// not pointers.<br>
- CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL);<br>
+ CurMapTypes.push_back(OMP_MAP_PRIVATE_VAL);<br>
CurSizes.push_back(CGF.getTypeSize(RI.getType()));<br>
} else {<br>
// Pointers are implicitly mapped with a zero size and no flags<br>
@@ -5539,9 +5652,8 @@ public:<br>
// default the value doesn't have to be retrieved. For an aggregate<br>
// type, the default is 'tofrom'.<br>
CurMapTypes.push_back(ElementType->isAggregateType()<br>
- ? (MappableExprsHandler::OMP_MAP_TO |<br>
- MappableExprsHandler::OMP_MAP_FROM)<br>
- : MappableExprsHandler::OMP_MAP_TO);<br>
+ ? (OMP_MAP_TO | OMP_MAP_FROM)<br>
+ : OMP_MAP_TO);<br>
<br>
// If we have a capture by reference we may need to add the private<br>
// pointer flag if the base declaration shows in some first-private<br>
@@ -5551,7 +5663,7 @@ public:<br>
}<br>
// Every default map produces a single argument, so, it is always the<br>
// first one.<br>
- CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF;<br>
+ CurMapTypes.back() |= OMP_MAP_FIRST_REF;<br>
}<br>
};<br>
<br>
@@ -5566,19 +5678,20 @@ enum OpenMPOffloadingReservedDeviceIDs {<br>
/// offloading runtime library. If there is no map or capture information,<br>
/// return nullptr by reference.<br>
static void<br>
-emitOffloadingArrays(CodeGenFunction &CGF, llvm::Value *&BasePointersArray,<br>
- llvm::Value *&PointersArray, llvm::Value *&SizesArray,<br>
- llvm::Value *&MapTypesArray,<br>
- MappableExprsHandler::MapValuesArrayTy &BasePointers,<br>
+emitOffloadingArrays(CodeGenFunction &CGF,<br>
+ MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,<br>
MappableExprsHandler::MapValuesArrayTy &Pointers,<br>
MappableExprsHandler::MapValuesArrayTy &Sizes,<br>
- MappableExprsHandler::MapFlagsArrayTy &MapTypes) {<br>
+ MappableExprsHandler::MapFlagsArrayTy &MapTypes,<br>
+ CGOpenMPRuntime::TargetDataInfo &Info) {<br>
auto &CGM = CGF.CGM;<br>
auto &Ctx = CGF.getContext();<br>
<br>
- BasePointersArray = PointersArray = SizesArray = MapTypesArray = nullptr;<br>
+ // Reset the array information.<br>
+ Info.clearArrayInfo();<br>
+ Info.NumberOfPtrs = BasePointers.size();<br>
<br>
- if (unsigned PointerNumVal = BasePointers.size()) {<br>
+ if (Info.NumberOfPtrs) {<br>
// Detect if we have any capture size requiring runtime evaluation of the<br>
// size so that a constant array could be eventually used.<br>
bool hasRuntimeEvaluationCaptureSize = false;<br>
@@ -5588,14 +5701,14 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
break;<br>
}<br>
<br>
- llvm::APInt PointerNumAP(32, PointerNumVal, /*isSigned=*/true);<br>
+ llvm::APInt PointerNumAP(32, Info.NumberOfPtrs, /*isSigned=*/true);<br>
QualType PointerArrayType =<br>
Ctx.getConstantArrayType(Ctx.VoidPtrTy, PointerNumAP, ArrayType::Normal,<br>
/*IndexTypeQuals=*/0);<br>
<br>
- BasePointersArray =<br>
+ Info.BasePointersArray =<br>
CGF.CreateMemTemp(PointerArrayType, ".offload_baseptrs").getPointer();<br>
- PointersArray =<br>
+ Info.PointersArray =<br>
CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs").getPointer();<br>
<br>
// If we don't have any VLA types or other types that require runtime<br>
@@ -5605,7 +5718,7 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
QualType SizeArrayType = Ctx.getConstantArrayType(<br>
Ctx.getSizeType(), PointerNumAP, ArrayType::Normal,<br>
/*IndexTypeQuals=*/0);<br>
- SizesArray =<br>
+ Info.SizesArray =<br>
CGF.CreateMemTemp(SizeArrayType, ".offload_sizes").getPointer();<br>
} else {<br>
// We expect all the sizes to be constant, so we collect them to create<br>
@@ -5621,7 +5734,7 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
/*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,<br>
SizesArrayInit, ".offload_sizes");<br>
SizesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);<br>
- SizesArray = SizesArrayGbl;<br>
+ Info.SizesArray = SizesArrayGbl;<br>
}<br>
<br>
// The map types are always constant so we don't need to generate code to<br>
@@ -5633,10 +5746,10 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
/*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,<br>
MapTypesArrayInit, ".offload_maptypes");<br>
MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);<br>
- MapTypesArray = MapTypesArrayGbl;<br>
+ Info.MapTypesArray = MapTypesArrayGbl;<br>
<br>
- for (unsigned i = 0; i < PointerNumVal; ++i) {<br>
- llvm::Value *BPVal = BasePointers[i];<br>
+ for (unsigned i = 0; i < Info.NumberOfPtrs; ++i) {<br>
+ llvm::Value *BPVal = *BasePointers[i];<br>
if (BPVal->getType()->isPointerTy())<br>
BPVal = CGF.Builder.CreateBitCast(BPVal, CGM.VoidPtrTy);<br>
else {<br>
@@ -5645,11 +5758,15 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
BPVal = CGF.Builder.CreateIntToPtr(BPVal, CGM.VoidPtrTy);<br>
}<br>
llvm::Value *BP = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), BasePointersArray,<br>
- 0, i);<br>
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),<br>
+ Info.BasePointersArray, 0, i);<br>
Address BPAddr(BP, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));<br>
CGF.Builder.CreateStore(BPVal, BPAddr);<br>
<br>
+ if (Info.requiresDevicePointerInfo())<br>
+ if (auto *DevVD = BasePointers[i].getDevicePtrDecl())<br>
+ Info.CaptureDeviceAddrMap.insert(std::make_pair(DevVD, BPAddr));<br>
+<br>
llvm::Value *PVal = Pointers[i];<br>
if (PVal->getType()->isPointerTy())<br>
PVal = CGF.Builder.CreateBitCast(PVal, CGM.VoidPtrTy);<br>
@@ -5659,14 +5776,15 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
PVal = CGF.Builder.CreateIntToPtr(PVal, CGM.VoidPtrTy);<br>
}<br>
llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), PointersArray, 0,<br>
- i);<br>
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),<br>
+ Info.PointersArray, 0, i);<br>
Address PAddr(P, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));<br>
CGF.Builder.CreateStore(PVal, PAddr);<br>
<br>
if (hasRuntimeEvaluationCaptureSize) {<br>
llvm::Value *S = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.SizeTy, PointerNumVal), SizesArray,<br>
+ llvm::ArrayType::get(CGM.SizeTy, Info.NumberOfPtrs),<br>
+ Info.SizesArray,<br>
/*Idx0=*/0,<br>
/*Idx1=*/i);<br>
Address SAddr(S, Ctx.getTypeAlignInChars(Ctx.getSizeType()));<br>
@@ -5682,23 +5800,24 @@ emitOffloadingArrays(CodeGenFunction &CG<br>
static void emitOffloadingArraysArgument(<br>
CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,<br>
llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,<br>
- llvm::Value *&MapTypesArrayArg, llvm::Value *BasePointersArray,<br>
- llvm::Value *PointersArray, llvm::Value *SizesArray,<br>
- llvm::Value *MapTypesArray, unsigned NumElems) {<br>
+ llvm::Value *&MapTypesArrayArg, CGOpenMPRuntime::TargetDataInfo &Info) {<br>
auto &CGM = CGF.CGM;<br>
- if (NumElems) {<br>
+ if (Info.NumberOfPtrs) {<br>
BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.VoidPtrTy, NumElems), BasePointersArray,<br>
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),<br>
+ Info.BasePointersArray,<br>
/*Idx0=*/0, /*Idx1=*/0);<br>
PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.VoidPtrTy, NumElems), PointersArray,<br>
+ llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),<br>
+ Info.PointersArray,<br>
/*Idx0=*/0,<br>
/*Idx1=*/0);<br>
SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.SizeTy, NumElems), SizesArray,<br>
+ llvm::ArrayType::get(CGM.SizeTy, Info.NumberOfPtrs), Info.SizesArray,<br>
/*Idx0=*/0, /*Idx1=*/0);<br>
MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(<br>
- llvm::ArrayType::get(CGM.Int32Ty, NumElems), MapTypesArray,<br>
+ llvm::ArrayType::get(CGM.Int32Ty, Info.NumberOfPtrs),<br>
+ Info.MapTypesArray,<br>
/*Idx0=*/0,<br>
/*Idx1=*/0);<br>
} else {<br>
@@ -5725,12 +5844,12 @@ void CGOpenMPRuntime::emitTargetCall(Cod<br>
<br>
// Fill up the arrays with all the captured variables.<br>
MappableExprsHandler::MapValuesArrayTy KernelArgs;<br>
- MappableExprsHandler::MapValuesArrayTy BasePointers;<br>
+ MappableExprsHandler::MapBaseValuesArrayTy BasePointers;<br>
MappableExprsHandler::MapValuesArrayTy Pointers;<br>
MappableExprsHandler::MapValuesArrayTy Sizes;<br>
MappableExprsHandler::MapFlagsArrayTy MapTypes;<br>
<br>
- MappableExprsHandler::MapValuesArrayTy CurBasePointers;<br>
+ MappableExprsHandler::MapBaseValuesArrayTy CurBasePointers;<br>
MappableExprsHandler::MapValuesArrayTy CurPointers;<br>
MappableExprsHandler::MapValuesArrayTy CurSizes;<br>
MappableExprsHandler::MapFlagsArrayTy CurMapTypes;<br>
@@ -5779,7 +5898,7 @@ void CGOpenMPRuntime::emitTargetCall(Cod<br>
<br>
// The kernel args are always the first elements of the base pointers<br>
// associated with a capture.<br>
- KernelArgs.push_back(CurBasePointers.front());<br>
+ KernelArgs.push_back(*CurBasePointers.front());<br>
// We need to append the results of this capture to what we already have.<br>
BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());<br>
Pointers.append(CurPointers.begin(), CurPointers.end());<br>
@@ -5802,17 +5921,11 @@ void CGOpenMPRuntime::emitTargetCall(Cod<br>
&D](CodeGenFunction &CGF, PrePostActionTy &) {<br>
auto &RT = CGF.CGM.getOpenMPRuntime();<br>
// Emit the offloading arrays.<br>
- llvm::Value *BasePointersArray;<br>
- llvm::Value *PointersArray;<br>
- llvm::Value *SizesArray;<br>
- llvm::Value *MapTypesArray;<br>
- emitOffloadingArrays(CGF, BasePointersArray, PointersArray, SizesArray,<br>
- MapTypesArray, BasePointers, Pointers, Sizes,<br>
- MapTypes);<br>
- emitOffloadingArraysArgument(CGF, BasePointersArray, PointersArray,<br>
- SizesArray, MapTypesArray, BasePointersArray,<br>
- PointersArray, SizesArray, MapTypesArray,<br>
- BasePointers.size());<br>
+ TargetDataInfo Info;<br>
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);<br>
+ emitOffloadingArraysArgument(CGF, Info.BasePointersArray,<br>
+ Info.PointersArray, Info.SizesArray,<br>
+ Info.MapTypesArray, Info);<br>
<br>
// On top of the arrays that were filled up, the target offloading call<br>
// takes as arguments the device id as well as the host pointer. The host<br>
@@ -5853,15 +5966,19 @@ void CGOpenMPRuntime::emitTargetCall(Cod<br>
assert(ThreadLimit && "Thread limit expression should be available along "<br>
"with number of teams.");<br>
llvm::Value *OffloadingArgs[] = {<br>
- DeviceID, OutlinedFnID, PointerNum,<br>
- BasePointersArray, PointersArray, SizesArray,<br>
- MapTypesArray, NumTeams, ThreadLimit};<br>
+ DeviceID, OutlinedFnID,<br>
+ PointerNum, Info.BasePointersArray,<br>
+ Info.PointersArray, Info.SizesArray,<br>
+ Info.MapTypesArray, NumTeams,<br>
+ ThreadLimit};<br>
Return = CGF.EmitRuntimeCall(<br>
RT.createRuntimeFunction(OMPRTL__tgt_target_teams), OffloadingArgs);<br>
} else {<br>
llvm::Value *OffloadingArgs[] = {<br>
- DeviceID, OutlinedFnID, PointerNum, BasePointersArray,<br>
- PointersArray, SizesArray, MapTypesArray};<br>
+ DeviceID, OutlinedFnID,<br>
+ PointerNum, Info.BasePointersArray,<br>
+ Info.PointersArray, Info.SizesArray,<br>
+ Info.MapTypesArray};<br>
Return = CGF.EmitRuntimeCall(RT.createRuntimeFunction(OMPRTL__tgt_target),<br>
OffloadingArgs);<br>
}<br>
@@ -6073,29 +6190,23 @@ void CGOpenMPRuntime::emitNumTeamsClause<br>
PushNumTeamsArgs);<br>
}<br>
<br>
-void CGOpenMPRuntime::emitTargetDataCalls(CodeGenFunction &CGF,<br>
- const OMPExecutableDirective &D,<br>
- const Expr *IfCond,<br>
- const Expr *Device,<br>
- const RegionCodeGenTy &CodeGen) {<br>
-<br>
+void CGOpenMPRuntime::emitTargetDataCalls(<br>
+ CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *IfCond,<br>
+ const Expr *Device, const RegionCodeGenTy &CodeGen, TargetDataInfo &Info) {<br>
if (!CGF.HaveInsertPoint())<br>
return;<br>
<br>
- llvm::Value *BasePointersArray = nullptr;<br>
- llvm::Value *PointersArray = nullptr;<br>
- llvm::Value *SizesArray = nullptr;<br>
- llvm::Value *MapTypesArray = nullptr;<br>
- unsigned NumOfPtrs = 0;<br>
+ // Action used to replace the default codegen action and turn privatization<br>
+ // off.<br>
+ PrePostActionTy NoPrivAction;<br>
<br>
// Generate the code for the opening of the data environment. Capture all the<br>
// arguments of the runtime call by reference because they are used in the<br>
// closing of the region.<br>
- auto &&BeginThenGen = [&D, &CGF, &BasePointersArray, &PointersArray,<br>
- &SizesArray, &MapTypesArray, Device,<br>
- &NumOfPtrs](CodeGenFunction &CGF, PrePostActionTy &) {<br>
+ auto &&BeginThenGen = [&D, &CGF, Device, &Info, &CodeGen, &NoPrivAction](<br>
+ CodeGenFunction &CGF, PrePostActionTy &) {<br>
// Fill up the arrays with all the mapped variables.<br>
- MappableExprsHandler::MapValuesArrayTy BasePointers;<br>
+ MappableExprsHandler::MapBaseValuesArrayTy BasePointers;<br>
MappableExprsHandler::MapValuesArrayTy Pointers;<br>
MappableExprsHandler::MapValuesArrayTy Sizes;<br>
MappableExprsHandler::MapFlagsArrayTy MapTypes;<br>
@@ -6103,21 +6214,16 @@ void CGOpenMPRuntime::emitTargetDataCall<br>
// Get map clause information.<br>
MappableExprsHandler MCHandler(D, CGF);<br>
MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);<br>
- NumOfPtrs = BasePointers.size();<br>
<br>
// Fill up the arrays and create the arguments.<br>
- emitOffloadingArrays(CGF, BasePointersArray, PointersArray, SizesArray,<br>
- MapTypesArray, BasePointers, Pointers, Sizes,<br>
- MapTypes);<br>
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);<br>
<br>
llvm::Value *BasePointersArrayArg = nullptr;<br>
llvm::Value *PointersArrayArg = nullptr;<br>
llvm::Value *SizesArrayArg = nullptr;<br>
llvm::Value *MapTypesArrayArg = nullptr;<br>
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,<br>
- SizesArrayArg, MapTypesArrayArg,<br>
- BasePointersArray, PointersArray, SizesArray,<br>
- MapTypesArray, NumOfPtrs);<br>
+ SizesArrayArg, MapTypesArrayArg, Info);<br>
<br>
// Emit device ID if any.<br>
llvm::Value *DeviceID = nullptr;<br>
@@ -6128,7 +6234,7 @@ void CGOpenMPRuntime::emitTargetDataCall<br>
DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF);<br>
<br>
// Emit the number of elements in the offloading arrays.<br>
- auto *PointerNum = CGF.Builder.getInt32(NumOfPtrs);<br>
+ auto *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);<br>
<br>
llvm::Value *OffloadingArgs[] = {<br>
DeviceID, PointerNum, BasePointersArrayArg,<br>
@@ -6136,23 +6242,24 @@ void CGOpenMPRuntime::emitTargetDataCall<br>
auto &RT = CGF.CGM.getOpenMPRuntime();<br>
CGF.EmitRuntimeCall(RT.createRuntimeFunction(OMPRTL__tgt_target_data_begin),<br>
OffloadingArgs);<br>
+<br>
+ // If device pointer privatization is required, emit the body of the region<br>
+ // here. It will have to be duplicated: with and without privatization.<br>
+ if (!Info.CaptureDeviceAddrMap.empty())<br>
+ CodeGen(CGF);<br>
};<br>
<br>
// Generate code for the closing of the data region.<br>
- auto &&EndThenGen = [&CGF, &BasePointersArray, &PointersArray, &SizesArray,<br>
- &MapTypesArray, Device,<br>
- &NumOfPtrs](CodeGenFunction &CGF, PrePostActionTy &) {<br>
- assert(BasePointersArray && PointersArray && SizesArray && MapTypesArray &&<br>
- NumOfPtrs && "Invalid data environment closing arguments.");<br>
+ auto &&EndThenGen = [&CGF, Device, &Info](CodeGenFunction &CGF,<br>
+ PrePostActionTy &) {<br>
+ assert(Info.isValid() && "Invalid data environment closing arguments.");<br>
<br>
llvm::Value *BasePointersArrayArg = nullptr;<br>
llvm::Value *PointersArrayArg = nullptr;<br>
llvm::Value *SizesArrayArg = nullptr;<br>
llvm::Value *MapTypesArrayArg = nullptr;<br>
emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,<br>
- SizesArrayArg, MapTypesArrayArg,<br>
- BasePointersArray, PointersArray, SizesArray,<br>
- MapTypesArray, NumOfPtrs);<br>
+ SizesArrayArg, MapTypesArrayArg, Info);<br>
<br>
// Emit device ID if any.<br>
llvm::Value *DeviceID = nullptr;<br>
@@ -6163,7 +6270,7 @@ void CGOpenMPRuntime::emitTargetDataCall<br>
DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF);<br>
<br>
// Emit the number of elements in the offloading arrays.<br>
- auto *PointerNum = CGF.Builder.getInt32(NumOfPtrs);<br>
+ auto *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);<br>
<br>
llvm::Value *OffloadingArgs[] = {<br>
DeviceID, PointerNum, BasePointersArrayArg,<br>
@@ -6173,24 +6280,40 @@ void CGOpenMPRuntime::emitTargetDataCall<br>
OffloadingArgs);<br>
};<br>
<br>
- // In the event we get an if clause, we don't have to take any action on the<br>
- // else side.<br>
- auto &&ElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {};<br>
+ // If we need device pointer privatization, we need to emit the body of the<br>
+ // region with no privatization in the 'else' branch of the conditional.<br>
+ // Otherwise, we don't have to do anything.<br>
+ auto &&BeginElseGen = [&Info, &CodeGen, &NoPrivAction](CodeGenFunction &CGF,<br>
+ PrePostActionTy &) {<br>
+ if (!Info.CaptureDeviceAddrMap.empty()) {<br>
+ CodeGen.setAction(NoPrivAction);<br>
+ CodeGen(CGF);<br>
+ }<br>
+ };<br>
+<br>
+ // We don't have to do anything to close the region if the if clause evaluates<br>
+ // to false.<br>
+ auto &&EndElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {};<br>
<br>
if (IfCond) {<br>
- emitOMPIfClause(CGF, IfCond, BeginThenGen, ElseGen);<br>
+ emitOMPIfClause(CGF, IfCond, BeginThenGen, BeginElseGen);<br>
} else {<br>
- RegionCodeGenTy BeginThenRCG(BeginThenGen);<br>
- BeginThenRCG(CGF);<br>
+ RegionCodeGenTy RCG(BeginThenGen);<br>
+ RCG(CGF);<br>
}<br>
<br>
- CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_target_data, CodeGen);<br>
+ // If we don't require privatization of device pointers, we emit the body in<br>
+ // between the runtime calls. This avoids duplicating the body code.<br>
+ if (Info.CaptureDeviceAddrMap.empty()) {<br>
+ CodeGen.setAction(NoPrivAction);<br>
+ CodeGen(CGF);<br>
+ }<br>
<br>
if (IfCond) {<br>
- emitOMPIfClause(CGF, IfCond, EndThenGen, ElseGen);<br>
+ emitOMPIfClause(CGF, IfCond, EndThenGen, EndElseGen);<br>
} else {<br>
- RegionCodeGenTy EndThenRCG(EndThenGen);<br>
- EndThenRCG(CGF);<br>
+ RegionCodeGenTy RCG(EndThenGen);<br>
+ RCG(CGF);<br>
}<br>
}<br>
<br>
@@ -6208,7 +6331,7 @@ void CGOpenMPRuntime::emitTargetDataStan<br>
// Generate the code for the opening of the data environment.<br>
auto &&ThenGen = [&D, &CGF, Device](CodeGenFunction &CGF, PrePostActionTy &) {<br>
// Fill up the arrays with all the mapped variables.<br>
- MappableExprsHandler::MapValuesArrayTy BasePointers;<br>
+ MappableExprsHandler::MapBaseValuesArrayTy BasePointers;<br>
MappableExprsHandler::MapValuesArrayTy Pointers;<br>
MappableExprsHandler::MapValuesArrayTy Sizes;<br>
MappableExprsHandler::MapFlagsArrayTy MapTypes;<br>
@@ -6217,19 +6340,12 @@ void CGOpenMPRuntime::emitTargetDataStan<br>
MappableExprsHandler MEHandler(D, CGF);<br>
MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);<br>
<br>
- llvm::Value *BasePointersArrayArg = nullptr;<br>
- llvm::Value *PointersArrayArg = nullptr;<br>
- llvm::Value *SizesArrayArg = nullptr;<br>
- llvm::Value *MapTypesArrayArg = nullptr;<br>
-<br>
// Fill up the arrays and create the arguments.<br>
- emitOffloadingArrays(CGF, BasePointersArrayArg, PointersArrayArg,<br>
- SizesArrayArg, MapTypesArrayArg, BasePointers,<br>
- Pointers, Sizes, MapTypes);<br>
- emitOffloadingArraysArgument(<br>
- CGF, BasePointersArrayArg, PointersArrayArg, SizesArrayArg,<br>
- MapTypesArrayArg, BasePointersArrayArg, PointersArrayArg, SizesArrayArg,<br>
- MapTypesArrayArg, BasePointers.size());<br>
+ TargetDataInfo Info;<br>
+ emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);<br>
+ emitOffloadingArraysArgument(CGF, Info.BasePointersArray,<br>
+ Info.PointersArray, Info.SizesArray,<br>
+ Info.MapTypesArray, Info);<br>
<br>
// Emit device ID if any.<br>
llvm::Value *DeviceID = nullptr;<br>
@@ -6243,8 +6359,8 @@ void CGOpenMPRuntime::emitTargetDataStan<br>
auto *PointerNum = CGF.Builder.getInt32(BasePointers.size());<br>
<br>
llvm::Value *OffloadingArgs[] = {<br>
- DeviceID, PointerNum, BasePointersArrayArg,<br>
- PointersArrayArg, SizesArrayArg, MapTypesArrayArg};<br>
+ DeviceID, PointerNum, Info.BasePointersArray,<br>
+ Info.PointersArray, Info.SizesArray, Info.MapTypesArray};<br>
<br>
auto &RT = CGF.CGM.getOpenMPRuntime();<br>
// Select the right runtime function call for each expected standalone<br>
<br>
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original)<br>
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Thu Jul 28 09:23:26 2016<br>
@@ -997,17 +997,59 @@ public:<br>
virtual void emitNumTeamsClause(CodeGenFunction &CGF, const Expr *NumTeams,<br>
const Expr *ThreadLimit, SourceLocation Loc);<br>
<br>
+ /// Struct that keeps all the relevant information that should be kept<br>
+ /// throughout a 'target data' region.<br>
+ class TargetDataInfo {<br>
+ /// Set to true if device pointer information have to be obtained.<br>
+ bool RequiresDevicePointerInfo = false;<br>
+<br>
+ public:<br>
+ /// The array of base pointer passed to the runtime library.<br>
+ llvm::Value *BasePointersArray = nullptr;<br>
+ /// The array of section pointers passed to the runtime library.<br>
+ llvm::Value *PointersArray = nullptr;<br>
+ /// The array of sizes passed to the runtime library.<br>
+ llvm::Value *SizesArray = nullptr;<br>
+ /// The array of map types passed to the runtime library.<br>
+ llvm::Value *MapTypesArray = nullptr;<br>
+ /// The total number of pointers passed to the runtime library.<br>
+ unsigned NumberOfPtrs = 0u;<br>
+ /// Map between the a declaration of a capture and the corresponding base<br>
+ /// pointer address where the runtime returns the device pointers.<br>
+ llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;<br>
+<br>
+ explicit TargetDataInfo() {}<br>
+ explicit TargetDataInfo(bool RequiresDevicePointerInfo)<br>
+ : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {}<br>
+ /// Clear information about the data arrays.<br>
+ void clearArrayInfo() {<br>
+ BasePointersArray = nullptr;<br>
+ PointersArray = nullptr;<br>
+ SizesArray = nullptr;<br>
+ MapTypesArray = nullptr;<br>
+ NumberOfPtrs = 0u;<br>
+ }<br>
+ /// Return true if the current target data information has valid arrays.<br>
+ bool isValid() {<br>
+ return BasePointersArray && PointersArray && SizesArray &&<br>
+ MapTypesArray && NumberOfPtrs;<br>
+ }<br>
+ bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }<br>
+ };<br>
+<br>
/// \brief Emit the target data mapping code associated with \a D.<br>
/// \param D Directive to emit.<br>
- /// \param IfCond Expression evaluated in if clause associated with the target<br>
- /// directive, or null if no if clause is used.<br>
+ /// \param IfCond Expression evaluated in if clause associated with the<br>
+ /// target directive, or null if no device clause is used.<br>
/// \param Device Expression evaluated in device clause associated with the<br>
/// target directive, or null if no device clause is used.<br>
- /// \param CodeGen Function that emits the enclosed region.<br>
+ /// \param Info A record used to store information that needs to be preserved<br>
+ /// until the region is closed.<br>
virtual void emitTargetDataCalls(CodeGenFunction &CGF,<br>
const OMPExecutableDirective &D,<br>
const Expr *IfCond, const Expr *Device,<br>
- const RegionCodeGenTy &CodeGen);<br>
+ const RegionCodeGenTy &CodeGen,<br>
+ TargetDataInfo &Info);<br>
<br>
/// \brief Emit the data mapping/movement code associated with the directive<br>
/// \a D that should be of the form 'target [{enter|exit} data | update]'.<br>
<br>
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)<br>
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Thu Jul 28 09:23:26 2016<br>
@@ -3400,22 +3400,137 @@ CodeGenFunction::getOMPCancelDestination<br>
return BreakContinueStack.back().BreakBlock;<br>
}<br>
<br>
+void CodeGenFunction::EmitOMPUseDevicePtrClause(<br>
+ const OMPClause &NC, OMPPrivateScope &PrivateScope,<br>
+ const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap) {<br>
+ const auto &C = cast<OMPUseDevicePtrClause>(NC);<br>
+ auto OrigVarIt = C.varlist_begin();<br>
+ auto InitIt = C.inits().begin();<br>
+ for (auto PvtVarIt : C.private_copies()) {<br>
+ auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(*OrigVarIt)->getDecl());<br>
+ auto *InitVD = cast<VarDecl>(cast<DeclRefExpr>(*InitIt)->getDecl());<br>
+ auto *PvtVD = cast<VarDecl>(cast<DeclRefExpr>(PvtVarIt)->getDecl());<br>
+<br>
+ // In order to identify the right initializer we need to match the<br>
+ // declaration used by the mapping logic. In some cases we may get<br>
+ // OMPCapturedExprDecl that refers to the original declaration.<br>
+ const ValueDecl *MatchingVD = OrigVD;<br>
+ if (auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) {<br>
+ // OMPCapturedExprDecl are used to privative fields of the current<br>
+ // structure.<br>
+ auto *ME = cast<MemberExpr>(OED->getInit());<br>
+ assert(isa<CXXThisExpr>(ME->getBase()) &&<br>
+ "Base should be the current struct!");<br>
+ MatchingVD = ME->getMemberDecl();<br>
+ }<br>
+<br>
+ // If we don't have information about the current list item, move on to<br>
+ // the next one.<br>
+ auto InitAddrIt = CaptureDeviceAddrMap.find(MatchingVD);<br>
+ if (InitAddrIt == CaptureDeviceAddrMap.end())<br>
+ continue;<br>
+<br>
+ bool IsRegistered = PrivateScope.addPrivate(OrigVD, [&]() -> Address {<br>
+ // Initialize the temporary initialization variable with the address we<br>
+ // get from the runtime library. We have to cast the source address<br>
+ // because it is always a void *. References are materialized in the<br>
+ // privatization scope, so the initialization here disregards the fact<br>
+ // the original variable is a reference.<br>
+ QualType AddrQTy =<br>
+ getContext().getPointerType(OrigVD->getType().getNonReferenceType());<br>
+ llvm::Type *AddrTy = ConvertTypeForMem(AddrQTy);<br>
+ Address InitAddr = Builder.CreateBitCast(InitAddrIt->second, AddrTy);<br>
+ setAddrOfLocalVar(InitVD, InitAddr);<br>
+<br>
+ // Emit private declaration, it will be initialized by the value we<br>
+ // declaration we just added to the local declarations map.<br>
+ EmitDecl(*PvtVD);<br>
+<br>
+ // The initialization variables reached its purpose in the emission<br>
+ // ofthe previous declaration, so we don't need it anymore.<br>
+ LocalDeclMap.erase(InitVD);<br>
+<br>
+ // Return the address of the private variable.<br>
+ return GetAddrOfLocalVar(PvtVD);<br>
+ });<br>
+ assert(IsRegistered && "firstprivate var already registered as private");<br>
+ // Silence the warning about unused variable.<br>
+ (void)IsRegistered;<br>
+<br>
+ ++OrigVarIt;<br>
+ ++InitIt;<br>
+ }<br>
+}<br>
+<br>
// Generate the instructions for '#pragma omp target data' directive.<br>
void CodeGenFunction::EmitOMPTargetDataDirective(<br>
const OMPTargetDataDirective &S) {<br>
- // The target data enclosed region is implemented just by emitting the<br>
- // statement.<br>
- auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) {<br>
- CGF.EmitStmt(cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());<br>
+ CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true);<br>
+<br>
+ // Create a pre/post action to signal the privatization of the device pointer.<br>
+ // This action can be replaced by the OpenMP runtime code generation to<br>
+ // deactivate privatization.<br>
+ bool PrivatizeDevicePointers = false;<br>
+ class DevicePointerPrivActionTy : public PrePostActionTy {<br>
+ bool &PrivatizeDevicePointers;<br>
+<br>
+ public:<br>
+ explicit DevicePointerPrivActionTy(bool &PrivatizeDevicePointers)<br>
+ : PrePostActionTy(), PrivatizeDevicePointers(PrivatizeDevicePointers) {}<br>
+ void Enter(CodeGenFunction &CGF) override {<br>
+ PrivatizeDevicePointers = true;<br>
+ }<br>
};<br>
+ DevicePointerPrivActionTy PrivAction(PrivatizeDevicePointers);<br>
+<br>
+ auto &&CodeGen = [&S, &Info, &PrivatizeDevicePointers](<br>
+ CodeGenFunction &CGF, PrePostActionTy &Action) {<br>
+ auto &&InnermostCodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) {<br>
+ CGF.EmitStmt(<br>
+ cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());<br>
+ };<br>
+<br>
+ // Codegen that selects wheather to generate the privatization code or not.<br>
+ auto &&PrivCodeGen = [&S, &Info, &PrivatizeDevicePointers,<br>
+ &InnermostCodeGen](CodeGenFunction &CGF,<br>
+ PrePostActionTy &Action) {<br>
+ RegionCodeGenTy RCG(InnermostCodeGen);<br>
+ PrivatizeDevicePointers = false;<br>
+<br>
+ // Call the pre-action to change the status of PrivatizeDevicePointers if<br>
+ // needed.<br>
+ Action.Enter(CGF);<br>
+<br>
+ if (PrivatizeDevicePointers) {<br>
+ OMPPrivateScope PrivateScope(CGF);<br>
+ // Emit all instances of the use_device_ptr clause.<br>
+ for (const auto *C : S.getClausesOfKind<OMPUseDevicePtrClause>())<br>
+ CGF.EmitOMPUseDevicePtrClause(*C, PrivateScope,<br>
+ Info.CaptureDeviceAddrMap);<br>
+ (void)PrivateScope.Privatize();<br>
+ RCG(CGF);<br>
+ } else<br>
+ RCG(CGF);<br>
+ };<br>
+<br>
+ // Forward the provided action to the privatization codegen.<br>
+ RegionCodeGenTy PrivRCG(PrivCodeGen);<br>
+ PrivRCG.setAction(Action);<br>
+<br>
+ // Notwithstanding the body of the region is emitted as inlined directive,<br>
+ // we don't use an inline scope as changes in the references inside the<br>
+ // region are expected to be visible outside, so we do not privative them.<br>
+ OMPLexicalScope Scope(CGF, S);<br>
+ CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_target_data,<br>
+ PrivRCG);<br>
+ };<br>
+<br>
+ RegionCodeGenTy RCG(CodeGen);<br>
<br>
// If we don't have target devices, don't bother emitting the data mapping<br>
// code.<br>
if (CGM.getLangOpts().OMPTargetTriples.empty()) {<br>
- OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);<br>
-<br>
- CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_target_data,<br>
- CodeGen);<br>
+ RCG(*this);<br>
return;<br>
}<br>
<br>
@@ -3429,7 +3544,12 @@ void CodeGenFunction::EmitOMPTargetDataD<br>
if (auto *C = S.getSingleClause<OMPDeviceClause>())<br>
Device = C->getDevice();<br>
<br>
- CGM.getOpenMPRuntime().emitTargetDataCalls(*this, S, IfCond, Device, CodeGen);<br>
+ // Set the action to signal privatization of device pointers.<br>
+ RCG.setAction(PrivAction);<br>
+<br>
+ // Emit region code.<br>
+ CGM.getOpenMPRuntime().emitTargetDataCalls(*this, S, IfCond, Device, RCG,<br>
+ Info);<br>
}<br>
<br>
void CodeGenFunction::EmitOMPTargetEnterDataDirective(<br>
<br>
Modified: cfe/trunk/lib/CodeGen/CodeGenFunction.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenFunction.h?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenFunction.h?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/CodeGen/CodeGenFunction.h (original)<br>
+++ cfe/trunk/lib/CodeGen/CodeGenFunction.h Thu Jul 28 09:23:26 2016<br>
@@ -2392,6 +2392,9 @@ public:<br>
OMPPrivateScope &PrivateScope);<br>
void EmitOMPPrivateClause(const OMPExecutableDirective &D,<br>
OMPPrivateScope &PrivateScope);<br>
+ void EmitOMPUseDevicePtrClause(<br>
+ const OMPClause &C, OMPPrivateScope &PrivateScope,<br>
+ const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap);<br>
/// \brief Emit code for copyin clause in \a D directive. The next code is<br>
/// generated at the start of outlined functions for directives:<br>
/// \code<br>
<br>
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)<br>
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Thu Jul 28 09:23:26 2016<br>
@@ -11800,7 +11800,10 @@ OMPClause *Sema::ActOnOpenMPUseDevicePtr<br>
SourceLocation StartLoc,<br>
SourceLocation LParenLoc,<br>
SourceLocation EndLoc) {<br>
- SmallVector<Expr *, 8> Vars;<br>
+ MappableVarListInfo MVLI(VarList);<br>
+ SmallVector<Expr *, 8> PrivateCopies;<br>
+ SmallVector<Expr *, 8> Inits;<br>
+<br>
for (auto &RefExpr : VarList) {<br>
assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause.");<br>
SourceLocation ELoc;<br>
@@ -11809,27 +11812,73 @@ OMPClause *Sema::ActOnOpenMPUseDevicePtr<br>
auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange);<br>
if (Res.second) {<br>
// It will be analyzed later.<br>
- Vars.push_back(RefExpr);<br>
+ MVLI.ProcessedVarList.push_back(RefExpr);<br>
+ PrivateCopies.push_back(nullptr);<br>
+ Inits.push_back(nullptr);<br>
}<br>
ValueDecl *D = Res.first;<br>
if (!D)<br>
continue;<br>
<br>
QualType Type = D->getType();<br>
- // item should be a pointer or reference to pointer<br>
- if (!Type.getNonReferenceType()->isPointerType()) {<br>
+ Type = Type.getNonReferenceType().getUnqualifiedType();<br>
+<br>
+ auto *VD = dyn_cast<VarDecl>(D);<br>
+<br>
+ // Item should be a pointer or reference to pointer.<br>
+ if (!Type->isPointerType()) {<br>
Diag(ELoc, diag::err_omp_usedeviceptr_not_a_pointer)<br>
<< 0 << RefExpr->getSourceRange();<br>
continue;<br>
}<br>
- Vars.push_back(RefExpr->IgnoreParens());<br>
+<br>
+ // Build the private variable and the expression that refers to it.<br>
+ auto VDPrivate = buildVarDecl(*this, ELoc, Type, D->getName(),<br>
+ D->hasAttrs() ? &D->getAttrs() : nullptr);<br>
+ if (VDPrivate->isInvalidDecl())<br>
+ continue;<br>
+<br>
+ CurContext->addDecl(VDPrivate);<br>
+ auto VDPrivateRefExpr = buildDeclRefExpr(<br>
+ *this, VDPrivate, RefExpr->getType().getUnqualifiedType(), ELoc);<br>
+<br>
+ // Add temporary variable to initialize the private copy of the pointer.<br>
+ auto *VDInit =<br>
+ buildVarDecl(*this, RefExpr->getExprLoc(), Type, ".devptr.temp");<br>
+ auto *VDInitRefExpr = buildDeclRefExpr(*this, VDInit, RefExpr->getType(),<br>
+ RefExpr->getExprLoc());<br>
+ AddInitializerToDecl(VDPrivate,<br>
+ DefaultLvalueConversion(VDInitRefExpr).get(),<br>
+ /*DirectInit=*/false, /*TypeMayContainAuto=*/false);<br>
+<br>
+ // If required, build a capture to implement the privatization initialized<br>
+ // with the current list item value.<br>
+ DeclRefExpr *Ref = nullptr;<br>
+ if (!VD)<br>
+ Ref = buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);<br>
+ MVLI.ProcessedVarList.push_back(VD ? RefExpr->IgnoreParens() : Ref);<br>
+ PrivateCopies.push_back(VDPrivateRefExpr);<br>
+ Inits.push_back(VDInitRefExpr);<br>
+<br>
+ // We need to add a data sharing attribute for this variable to make sure it<br>
+ // is correctly captured. A variable that shows up in a use_device_ptr has<br>
+ // similar properties of a first private variable.<br>
+ DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);<br>
+<br>
+ // Create a mappable component for the list item. List items in this clause<br>
+ // only need a component.<br>
+ MVLI.VarBaseDeclarations.push_back(D);<br>
+ MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);<br>
+ MVLI.VarComponents.back().push_back(<br>
+ OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));<br>
}<br>
<br>
- if (Vars.empty())<br>
+ if (MVLI.ProcessedVarList.empty())<br>
return nullptr;<br>
<br>
- return OMPUseDevicePtrClause::Create(Context, StartLoc, LParenLoc, EndLoc,<br>
- Vars);<br>
+ return OMPUseDevicePtrClause::Create(<br>
+ Context, StartLoc, LParenLoc, EndLoc, MVLI.ProcessedVarList,<br>
+ PrivateCopies, Inits, MVLI.VarBaseDeclarations, MVLI.VarComponents);<br>
}<br>
<br>
OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,<br>
<br>
Modified: cfe/trunk/lib/Serialization/ASTReaderStmt.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderStmt.cpp?rev=276977&r1=276976&r2=276977&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderStmt.cpp?rev=276977&r1=276976&r2=276977&view=diff</a><br>
==============================================================================<br>
--- cfe/trunk/lib/Serialization/ASTReaderStmt.cpp (original)<br>
+++ cfe/tr</blockquote></div><div dir="ltr">-- <br></div><div data-smartmail="gmail_signature">Mike<br>Sent from phone</div>