<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)">
<style><!--
/* Font Definitions */
@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;}
/* 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:#0563C1;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:#954F72;
text-decoration:underline;}
p.msonormal0, li.msonormal0, div.msonormal0
{mso-style-name:msonormal;
mso-margin-top-alt:auto;
margin-right:0in;
mso-margin-bottom-alt:auto;
margin-left:0in;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
span.EmailStyle18
{mso-style-type:personal;
font-family:"Calibri",sans-serif;
color:windowtext;}
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:56.7pt 42.5pt 56.7pt 85.05pt;}
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="#0563C1" vlink="#954F72">
<div class="WordSection1">
<p class="MsoNormal">Right. I am taking a look.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Thanks.<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> Bessonova, Kristina [mailto:kristina.bessonova@intel.com]
<br>
<b>Sent:</b> Monday, February 05, 2018 12:15 PM<br>
<b>To:</b> cfe-dev@lists.llvm.org<br>
<b>Cc:</b> Sumner, Brian <Brian.Sumner@amd.com>; Liu, Yaxun (Sam) <Yaxun.Liu@amd.com>; anastasia.stulova@arm.com<br>
<b>Subject:</b> [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>
<p class="MsoNormal">Hi all,<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">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></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">$ cat test.cl<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">void foo(size_t id, __global int* out) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""> out[id] = id;<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">}<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">kernel void enqueue_foo(__global int* out) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""> size_t id = get_global_id(0);<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""> void (^fooBlock)(void) = ^{ foo(id, out); };<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""> queue_t queue = get_default_queue();<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""> ndrange_t ndrange = ndrange_1D(1);<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""> enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, fooBlock);<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">}<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">$ build/bin/clang --version<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">clang version 7.0.0 (<a href="https://git.llvm.org/git/clang.git/">https://git.llvm.org/git/clang.git/</a> 08712fff7fba84b88e2e57b3c739d53b1aab1ed6)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">$ clang -cc1 -emit-llvm -x cl -triple spir64-unknown-unknown -finclude-default-header -cl-std=CL2.0 test.cl<o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">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*<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">>::ret_type = const clang::BlockExpr*]: Assertion `isa<X>(Val) && "cast<Ty>() argument of incompatible type!"' failed.<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">…<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#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<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">.h:257:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#10 clang::CodeGen::CGOpenCLRuntime::emitOpenCLEnqueuedBlock(clang::CodeGen::CodeGenFunction&, clang::Expr const*) ./llvm/tools/clang/lib/CodeGen/CGOpenCLRuntime.cpp:125:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#11 clang::CodeGen::CodeGenFunction::EmitBuiltinExpr(clang::FunctionDecl const*, unsigned int, clang::CallExpr const*, clang::CodeGen::ReturnValueSlot) ./llvm/tools/clang/<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">lib/CodeGen/CGBuiltin.cpp:3017:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#12 clang::CodeGen::CodeGenFunction::EmitCallExpr(clang::CallExpr const*, clang::CodeGen::ReturnValueSlot) ./llvm/tools/clang/lib/CodeGen/CGExpr.cpp:4218:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#13 (anonymous namespace)::ScalarExprEmitter::VisitCallExpr(clang::CallExpr const*) ./llvm/tools/clang/lib/CodeGen/CGExprScalar.cpp:517:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#14 clang::StmtVisitorBase<clang::make_ptr, (anonymous namespace)::ScalarExprEmitter, llvm::Value*>::Visit(clang::Stmt*) ./llvm/build/tools/clang/include/clang/AST/StmtNod<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">es.inc:329:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New"">#15 (anonymous namespace)::ScalarExprEmitter::Visit(clang::Expr*) ./llvm/tools/clang/lib/CodeGen/CGExprScalar.cpp:358:0<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal">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></p>
<p class="MsoNormal">So, it seems we have a bug here. Am I right?<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Thanks,<o:p></o:p></p>
<p class="MsoNormal">Kristina<o:p></o:p></p>
<p><span lang="RU"><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<o:p></o:p></span></p>
<p><span lang="RU">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.<o:p></o:p></span></p>
</div>
</body>
</html>