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