r276977 - [OpenMP] Codegen for use_device_ptr clause.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 29 22:01:11 PDT 2016
Hans, Samuel,
I'm ok with this fix since it is required for compatibility with MSVC.
Отправлено с iPhone
30 июля 2016 г., в 0:20, Samuel F Antao via cfe-commits <cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>> написал(а):
Hi Hans,
Thanks for the troubleshooting. I'd prefer to avoid the 'this->', but if Alexey is okay with that I wouldn't mind. The best solution would be to update the compiler of the bot if there is a newer version and see if the issue goes away, but not sure if that is possible at all.
Thanks again,
Samuel
----- Original message -----
From: Hans Wennborg <hans at chromium.org<mailto:hans at chromium.org>>
Sent by: hwennborg at google.com<mailto:hwennborg at google.com>
To: Samuel F Antao/Watson/IBM at IBMUS
Cc: Mike Aizatsky <aizatsky at google.com<mailto:aizatsky at google.com>>, cfe-commits <cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>>, "Bataev, Alexey" <a.bataev at hotmail.com<mailto:a.bataev at hotmail.com>>
Subject: Re: r276977 - [OpenMP] Codegen for use_device_ptr clause.
Date: Fri, Jul 29, 2016 5:04 PM
I agree this looks like a compiler bug, but I'm not sure what's
causing it exactly. Putting "this->" in front of those accesses to
CurDir and CGF seems to fix it, but it's also very ugly.
I tried a couple of other things, like moving the method out-of-line,
but no luck yet.
On Fri, Jul 29, 2016 at 1:45 PM, Samuel F Antao <sfantao at us.ibm.com<mailto:sfantao at us.ibm.com>> wrote:
> I don't have easy access to Windows machine to test this, but it seems that
> whichever compiler is being used is broken. This is about a member function
> accessing a variable of the same class. I don't think I am doing any
> non-portable code pattern here.
>
> Alexey, do you have a easy way to test this? Can you give me a little help
> here?
>
> Thanks!
> Samuel
>
>
> ----- Original message -----
> From: Hans Wennborg <hans at chromium.org<mailto:hans at chromium.org>>
> Sent by: hwennborg at google.com<mailto:hwennborg at google.com>
> To: Samuel F Antao/Watson/IBM at IBMUS
> Cc: Mike Aizatsky <aizatsky at google.com<mailto:aizatsky at google.com>>, cfe-commits
> <cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>>
> Subject: Re: r276977 - [OpenMP] Codegen for use_device_ptr clause.
> Date: Fri, Jul 29, 2016 3:53 PM
>
> Yes, it still appears to be broken:
>
> http://lab.llvm.org:8011/builders/sanitizer-windows/builds/26504/steps/run%20tests/logs/stdio
>
> On Thu, Jul 28, 2016 at 1:30 PM, Samuel F Antao via cfe-commits
> <cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>> wrote:
>> Hi Mike,
>>
>> I've already pushed r276981 and r276988 to fix those failures shortly
>> after
>> I pushed the patch. After that I didn't receive any bot message
>> complaining
>> about that code.
>>
>> Do you still see the issues?
>>
>> Thanks,
>> Samuel
>>
>>
>> ----- Original message -----
>> From: Mike Aizatsky <aizatsky at google.com<mailto:aizatsky at google.com>>
>> To: Samuel F Antao/Watson/IBM at IBMUS, cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>
>> Cc:
>> Subject: Re: r276977 - [OpenMP] Codegen for use_device_ptr clause.
>> Date: Thu, Jul 28, 2016 4:20 PM
>>
>> 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<mailto: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
>>
>>
>>
>>
>> _______________________________________________
>> cfe-commits mailing list
>> cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>
>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>>
>
>
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160730/cac21505/attachment-0001.html>
More information about the cfe-commits
mailing list