<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:0cm;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;
mso-fareast-language:EN-US;}
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;}
span.EmailStyle17
{mso-style-type:personal-compose;
font-family:"Calibri",sans-serif;
color:windowtext;}
.MsoChpDefault
{mso-style-type:export-only;
font-family:"Calibri",sans-serif;
mso-fareast-language:EN-US;}
@page WordSection1
{size:612.0pt 792.0pt;
margin:2.0cm 42.5pt 2.0cm 3.0cm;}
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="RU" link="#0563C1" vlink="#954F72">
<div class="WordSection1">
<p class="MsoNormal"><span lang="EN-US">Hi all,<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US"><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US">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="MsoNormal"><span lang="EN-US"><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New"">$ cat test.cl<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New"">void foo(size_t id, __global int* out) {<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""> out[id] = id;<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New"">}<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New"">kernel void enqueue_foo(__global int* out) {<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""> size_t id = get_global_id(0);<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""> void (^fooBlock)(void) = ^{ foo(id, out); };<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""> queue_t queue = get_default_queue();<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""> ndrange_t ndrange = ndrange_1D(1);<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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 lang="EN-US" style="font-family:"Courier New"">}<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New"">$ build/bin/clang --version<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New"">clang version 7.0.0 (https://git.llvm.org/git/clang.git/ 08712fff7fba84b88e2e57b3c739d53b1aab1ed6)<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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"><span lang="EN-US"><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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 lang="EN-US" 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 lang="EN-US" style="font-family:"Courier New"">…<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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 lang="EN-US" style="font-family:"Courier New"">.h:257:0<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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 lang="EN-US" 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 lang="EN-US" style="font-family:"Courier New"">lib/CodeGen/CGBuiltin.cpp:3017:0<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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 lang="EN-US" 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 lang="EN-US" 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 lang="EN-US" style="font-family:"Courier New"">es.inc:329:0<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US" 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 lang="EN-US" style="font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US">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="MsoNormal"><span lang="EN-US">So, it seems we have a bug here. Am I right?<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US"><o:p> </o:p></span></p>
<p class="MsoNormal"><span lang="EN-US">Thanks,<o:p></o:p></span></p>
<p class="MsoNormal"><span lang="EN-US">Kristina<o:p></o:p></span></p>
</div>
<p><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</p><p>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.</p>
</body>
</html>