r276977 - [OpenMP] Codegen for use_device_ptr clause.

Mike Aizatsky via cfe-commits cfe-commits at lists.llvm.org
Thu Jul 28 13:19:51 PDT 2016


Samuel,

I think your change breaks the build:

http://lab.llvm.org:8011/builders/sanitizer-windows/builds/26397

FAILED:
tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/CGOpenMPRuntime.cpp.obj
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
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
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
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488)
: error C2065: 'Directive' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488)
: error C2228: left of '.getClausesOfKind' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488)
: error C2059: syntax error : ')'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488)
: error C2143: syntax error : missing ';' before '<end Parse>'
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
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491)
: error C2065: 'Directive' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491)
: error C2228: left of '.getClausesOfKind' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491)
: error C2059: syntax error : ')'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491)
: error C2143: syntax error : missing ';' before '<end Parse>'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5493)
: error C2143: syntax error : missing ')' before ';'
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
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494)
: error C2065: 'Directive' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494)
: error C2228: left of '.getClausesOfKind' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494)
: error C2059: syntax error : ')'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494)
: error C2143: syntax error : missing ';' before '<end Parse>'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5496)
: error C2143: syntax error : missing ')' before ';'
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
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502)
: error C2065: 'Directive' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502)
: error C2228: left of '.getClausesOfKind' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502)
: error C2059: syntax error : ')'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502)
: error C2143: syntax error : missing ';' before '<end Parse>'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5503)
: error C2143: syntax error : missing ';' before '{'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5503)
: error C2143: syntax error : missing ')' before ';'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504)
: error C2065: 'L' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504)
: error C2228: left of '.second' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504)
: error C2228: left of '.empty' must have class/struct/union
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505)
: error C2065: 'L' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505)
: error C2228: left of '.second' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505)
: error C2228: left of '.back' must have class/struct/union
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505)
: error C2228: left of '.getAssociatedDeclaration' must have
class/struct/union
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507)
: error C2065: 'L' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507)
: error C2228: left of '.second' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507)
: error C2228: left of '.back' must have class/struct/union
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507)
: error C2228: left of '.getAssociatedExpression' must have
class/struct/union
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5511)
: error C3536: 'IE': cannot be used before it is initialized
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5523)
: error C3536: 'IE': cannot be used before it is initialized
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5528)
: error C2044: illegal continue
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
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535)
: error C2065: 'CGF' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535)
: error C2228: left of '.EmitLoadOfLValue' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535)
: error C2228: left of '.EmitLValue' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535)
: error C3536: 'IE': cannot be used before it is initialized
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5536)
: error C2228: left of '.getScalarVal' must have class/struct/union
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
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539)
: error C2065: 'CGF' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539)
: error C2228: left of '.SizeTy' must have class/struct/union
        type is 'unknown-type'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543)
: error C2143: syntax error : missing ')' before ';'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543)
: error C2143: syntax error : missing ';' before ')'
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543)
: error C2065: '$S239' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543)
: error C2065: '$S240' : undeclared identifier
C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543)
: error C2143: syntax error : missing ';' before '{'
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
18446743528493.273 [136/2/116] Building CXX object
tools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\CGOpenCLRuntime.cpp.obj

On Thu, Jul 28, 2016 at 7:31 AM Samuel Antao via cfe-commits <
cfe-commits at lists.llvm.org> wrote:

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

-- 
Mike
Sent from phone
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160728/8b69d3e5/attachment-0001.html>


More information about the cfe-commits mailing list