<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=us-ascii">
<meta name="Generator" content="Microsoft Word 15 (filtered medium)">
<!--[if !mso]><style>v\:* {behavior:url(#default#VML);}
o\:* {behavior:url(#default#VML);}
w\:* {behavior:url(#default#VML);}
.shape {behavior:url(#default#VML);}
</style><![endif]--><style><!--
/* Font Definitions */
@font-face
{font-family:SimSun;
panose-1:2 1 6 0 3 1 1 1 1 1;}
@font-face
{font-family:"Cambria Math";
panose-1:2 4 5 3 5 4 6 3 2 4;}
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
{font-family:"\@SimSun";
panose-1:2 1 6 0 3 1 1 1 1 1;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
p.msonormal0, li.msonormal0, div.msonormal0
{mso-style-name:msonormal;
margin:0in;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
p.xmsonormal, li.xmsonormal, div.xmsonormal
{mso-style-name:x_msonormal;
margin:0in;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
span.EmailStyle21
{mso-style-type:personal-reply;
font-family:"Calibri",sans-serif;
color:windowtext;}
.MsoChpDefault
{mso-style-type:export-only;
font-size:10.0pt;}
@page WordSection1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang="EN-US" link="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoNormal">Thanks Anastasia for investigating the issue.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">The wrapper kernel needs to call the block invoke function, which is created by
<span style="font-size:12.0pt;color:black">emitOpenCLEnqueuedBlock.</span><o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Sam<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<div>
<div style="border:none;border-top:solid #E1E1E1 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="MsoNormal"><b>From:</b> Anastasia Stulova [mailto:Anastasia.Stulova@arm.com]
<br>
<b>Sent:</b> Thursday, February 08, 2018 4:05 PM<br>
<b>To:</b> Liu, Yaxun (Sam) <Yaxun.Liu@amd.com>; Bessonova, Kristina <kristina.bessonova@intel.com>; cfe-dev@lists.llvm.org<br>
<b>Cc:</b> Sumner, Brian <Brian.Sumner@amd.com>; nd <nd@arm.com><br>
<b>Subject:</b> Re: [OpenCL] clang can't compile a simple enqueue_kernel with default opt level<o:p></o:p></p>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<div id="divtagdefaultwrapper">
<div id="divtagdefaultwrapper">
<p><span style="font-size:12.0pt;color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:12.0pt;color:black">It seems one of the assumptions that E in emitOpenCLEnqueuedBlock has BlockExpr type is not right and therefore the assert is triggered inside the static cast. This can be easily fixed by calling
IgnoreImplicit() method.<br>
<br>
But then there is another issue due to the generation of the Block during the AST parsing of the block declaration and during the enqueue_kernel builtin generation path. I think we should avoid generating the block literal from the enqueue_kernel generation.
This doesn't really match the original compilation flow.<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-size:12.0pt;color:black"><o:p> </o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-size:12.0pt;color:black">@Sam, I was just wondering whether we could avoid generating the literal inside the emitOpenCLEnqueuedBlock and pass the name of the block and num of its params to createEnqueuedBlockKernel or
perhaps we could pass the BlockExpr? The wrapper kernel function is pretty simple at the end and doesn't require the block invoke function itself (it can just rebuild the prototype). What do you think?<o:p></o:p></span></p>
</div>
<div>
<p class="MsoNormal"><span style="font-size:12.0pt;color:black"><o:p> </o:p></span></p>
</div>
<div>
<p class="MsoNormal" style="margin-bottom:12.0pt"><span style="font-size:12.0pt;color:black">Anastasia<o:p></o:p></span></p>
<div>
<div class="MsoNormal" align="center" style="text-align:center"><span style="font-size:12.0pt;color:black">
<hr size="2" width="98%" align="center">
</span></div>
<div id="divRplyFwdMsg">
<p class="MsoNormal"><b><span style="color:black">From:</span></b><span style="color:black"> Liu, Yaxun (Sam) <<a href="mailto:Yaxun.Liu@amd.com">Yaxun.Liu@amd.com</a>><br>
<b>Sent:</b> 05 February 2018 18:22<br>
<b>To:</b> Bessonova, Kristina; <a href="mailto:cfe-dev@lists.llvm.org">cfe-dev@lists.llvm.org</a><br>
<b>Cc:</b> Sumner, Brian; Anastasia Stulova<br>
<b>Subject:</b> RE: [OpenCL] clang can't compile a simple enqueue_kernel with default opt level</span><span style="font-size:12.0pt;color:black">
<o:p></o:p></span></p>
<div>
<p class="MsoNormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
</div>
</div>
<div>
<div>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Right. I am taking a look.<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Thanks.<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Sam<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<div>
<div style="border:none;border-top:solid #E1E1E1 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="xmsonormal"><b><span style="font-size:12.0pt;color:black">From:</span></b><span style="font-size:12.0pt;color:black"> Bessonova, Kristina [<a href="mailto:kristina.bessonova@intel.com">mailto:kristina.bessonova@intel.com</a>]
<br>
<b>Sent:</b> Monday, February 05, 2018 12:15 PM<br>
<b>To:</b> <a href="mailto:cfe-dev@lists.llvm.org">cfe-dev@lists.llvm.org</a><br>
<b>Cc:</b> Sumner, Brian <<a href="mailto:Brian.Sumner@amd.com">Brian.Sumner@amd.com</a>>; Liu, Yaxun (Sam) <<a href="mailto:Yaxun.Liu@amd.com">Yaxun.Liu@amd.com</a>>;
<a href="mailto:anastasia.stulova@arm.com">anastasia.stulova@arm.com</a><br>
<b>Subject:</b> [OpenCL] clang can't compile a simple enqueue_kernel with default opt level<o:p></o:p></span></p>
</div>
</div>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Hi all,<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">I’ve noticed that clang fails to compile a simple OpenCL kernel with an enqueue_kernel() call if optimization level isn’t set:<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">$ cat test.cl</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">void foo(size_t id, __global int* out) {</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> out[id] = id;</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">}</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> </span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">kernel void enqueue_foo(__global int* out) {</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> size_t id = get_global_id(0);</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> </span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> void (^fooBlock)(void) = ^{ foo(id, out); };</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> </span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> queue_t queue = get_default_queue();</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> ndrange_t ndrange = ndrange_1D(1);</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, fooBlock);</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">}</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> </span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">$ build/bin/clang --version</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">clang version 7.0.0 (<a href="https://git.llvm.org/git/clang.git/">https://git.llvm.org/git/clang.git/</a> 08712fff7fba84b88e2e57b3c739d53b1aab1ed6)</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> </span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">$ clang -cc1 -emit-llvm -x cl -triple spir64-unknown-unknown -finclude-default-header -cl-std=CL2.0 test.cl</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">clang: ./llvm/include/llvm/Support/Casting.h:255: typename llvm::cast_retty<X, Y*>::ret_type llvm::cast(Y*) [with X = clang::BlockExpr; Y = const clang::Expr; typename
llvm::cast_retty<X, Y*</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">>::ret_type = const clang::BlockExpr*]: Assertion `isa<X>(Val) && "cast<Ty>() argument of incompatible type!"' failed.</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">…</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#9 llvm::cast_retty<clang::BlockExpr, clang::Expr const*>::ret_type llvm::cast<clang::BlockExpr, clang::Expr const>(clang::Expr const*) ./llvm/include/llvm/Support/Casting</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">.h:257:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#10 clang::CodeGen::CGOpenCLRuntime::emitOpenCLEnqueuedBlock(clang::CodeGen::CodeGenFunction&, clang::Expr const*) ./llvm/tools/clang/lib/CodeGen/CGOpenCLRuntime.cpp:125:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#11 clang::CodeGen::CodeGenFunction::EmitBuiltinExpr(clang::FunctionDecl const*, unsigned int, clang::CallExpr const*, clang::CodeGen::ReturnValueSlot) ./llvm/tools/clang/</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">lib/CodeGen/CGBuiltin.cpp:3017:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#12 clang::CodeGen::CodeGenFunction::EmitCallExpr(clang::CallExpr const*, clang::CodeGen::ReturnValueSlot) ./llvm/tools/clang/lib/CodeGen/CGExpr.cpp:4218:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#13 (anonymous namespace)::ScalarExprEmitter::VisitCallExpr(clang::CallExpr const*) ./llvm/tools/clang/lib/CodeGen/CGExprScalar.cpp:517:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#14 clang::StmtVisitorBase<clang::make_ptr, (anonymous namespace)::ScalarExprEmitter, llvm::Value*>::Visit(clang::Stmt*) ./llvm/build/tools/clang/include/clang/AST/StmtNod</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">es.inc:329:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black">#15 (anonymous namespace)::ScalarExprEmitter::Visit(clang::Expr*) ./llvm/tools/clang/lib/CodeGen/CGExprScalar.cpp:358:0</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;font-family:"Courier New";color:black"> </span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Clang compiles OpenCL sources with ‘-O2’ optimization level by default. However tests for enqueue_kernel() in clang are compiled with ‘-O0’.<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">So, it seems we have a bug here. Am I right?<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black"> <o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Thanks,<o:p></o:p></span></p>
<p class="xmsonormal"><span style="font-size:12.0pt;color:black">Kristina<o:p></o:p></span></p>
<p><span lang="RU" style="font-size:12.0pt;color:black"><br>
--------------------------------------------------------------------<br>
Joint Stock Company Intel A/O<br>
Registered legal address: Krylatsky Hills Business Park, <br>
17 Krylatskaya Str., Bldg 4, Moscow 121614, <br>
Russian Federation</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
<p><span lang="RU" style="font-size:12.0pt;color:black">This e-mail and any attachments may contain confidential material for<br>
the sole use of the intended recipient(s). Any review or distribution<br>
by others is strictly prohibited. If you are not the intended<br>
recipient, please contact the sender and delete all copies.</span><span style="font-size:12.0pt;color:black"><o:p></o:p></span></p>
</div>
</div>
</div>
</div>
</div>
</div>
</body>
</html>