<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=Windows-1252">
<style type="text/css" style="display:none;"><!-- P {margin-top:0;margin-bottom:0;} --></style>
</head>
<body dir="ltr">
<div id="divtagdefaultwrapper" style="font-size:12pt;color:#000000;font-family:Calibri,Helvetica,sans-serif;" dir="ltr">
<div id="divtagdefaultwrapper" style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri,Helvetica,sans-serif,"EmojiFont","Apple Color Emoji","Segoe UI Emoji",NotoColorEmoji,"Segoe UI Symbol","Android Emoji",EmojiSymbols;" dir="ltr">
<p style="margin-top:0;margin-bottom:0"><br>
</p>
It seems one of the assumptions that E in <span>emitOpenCLEnqueuedBlock has BlockExpr type is not right and therefore the assert is triggered inside the static cast. This can be easily fixed by calling
<span>IgnoreImplicit</span>() 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.<br>
</span></div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri,Helvetica,sans-serif,"EmojiFont","Apple Color Emoji","Segoe UI Emoji",NotoColorEmoji,"Segoe UI Symbol","Android Emoji",EmojiSymbols;" dir="ltr">
<span><br>
</span></div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri,Helvetica,sans-serif,"EmojiFont","Apple Color Emoji","Segoe UI Emoji",NotoColorEmoji,"Segoe UI Symbol","Android Emoji",EmojiSymbols;" dir="ltr">
<span>@Sam, I was just wondering whether we could avoid generating the literal inside the
<span>emitOpenCLEnqueuedBlock</span> and pass the name of the block and num of its params to
<span>createEnqueuedBlockKernel</span> 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?</span></div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri,Helvetica,sans-serif,"EmojiFont","Apple Color Emoji","Segoe UI Emoji",NotoColorEmoji,"Segoe UI Symbol","Android Emoji",EmojiSymbols;" dir="ltr">
<span><br>
</span></div>
<div style="font-size: 12pt; color: rgb(0, 0, 0); font-family: Calibri,Helvetica,sans-serif,"EmojiFont","Apple Color Emoji","Segoe UI Emoji",NotoColorEmoji,"Segoe UI Symbol","Android Emoji",EmojiSymbols;" dir="ltr">
<span>Anastasia</span><br>
<br>
<div style="color: rgb(0, 0, 0);">
<hr style="display:inline-block;width:98%" tabindex="-1">
<div id="divRplyFwdMsg" dir="ltr"><font style="font-size:11pt" face="Calibri, sans-serif" color="#000000"><b>From:</b> Liu, Yaxun (Sam) <Yaxun.Liu@amd.com><br>
<b>Sent:</b> 05 February 2018 18:22<br>
<b>To:</b> Bessonova, Kristina; cfe-dev@lists.llvm.org<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</font>
<div> </div>
</div>
<div link="#0563C1" vlink="#954F72" lang="EN-US">
<div class="x_WordSection1">
<p class="x_MsoNormal">Right. I am taking a look.</p>
<p class="x_MsoNormal"> </p>
<p class="x_MsoNormal">Thanks.</p>
<p class="x_MsoNormal"> </p>
<p class="x_MsoNormal">Sam</p>
<p class="x_MsoNormal"> </p>
<div>
<div style="border:none; border-top:solid #E1E1E1 1.0pt; padding:3.0pt 0in 0in 0in">
<p class="x_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</p>
</div>
</div>
<p class="x_MsoNormal"> </p>
<p class="x_MsoNormal">Hi all,</p>
<p class="x_MsoNormal"> </p>
<p class="x_MsoNormal">I’ve noticed that clang fails to compile a simple OpenCL kernel with an enqueue_kernel() call if optimization level isn’t set:</p>
<p class="x_MsoNormal"> </p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">$ cat test.cl</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">void foo(size_t id, __global int* out) {</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> out[id] = id;</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">}</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> </span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">kernel void enqueue_foo(__global int* out) {</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> size_t id = get_global_id(0);</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> </span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> void (^fooBlock)(void) = ^{ foo(id, out); };</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> </span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> queue_t queue = get_default_queue();</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> ndrange_t ndrange = ndrange_1D(1);</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> enqueue_kernel(queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, fooBlock);</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">}</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> </span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">$ build/bin/clang --version</span></p>
<p class="x_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)</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> </span></p>
<p class="x_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</span></p>
<p class="x_MsoNormal"> </p>
<p class="x_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*</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">>::ret_type = const clang::BlockExpr*]: Assertion `isa<X>(Val) && "cast<Ty>() argument of incompatible type!"' failed.</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">…</span></p>
<p class="x_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</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">.h:257:0</span></p>
<p class="x_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</span></p>
<p class="x_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/</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">lib/CodeGen/CGBuiltin.cpp:3017:0</span></p>
<p class="x_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</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">#13 (anonymous namespace)::ScalarExprEmitter::VisitCallExpr(clang::CallExpr const*) ./llvm/tools/clang/lib/CodeGen/CGExprScalar.cpp:517:0</span></p>
<p class="x_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</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">es.inc:329:0</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New"">#15 (anonymous namespace)::ScalarExprEmitter::Visit(clang::Expr*) ./llvm/tools/clang/lib/CodeGen/CGExprScalar.cpp:358:0</span></p>
<p class="x_MsoNormal"><span style="font-family:"Courier New""> </span></p>
<p class="x_MsoNormal">Clang compiles OpenCL sources with ‘-O2’ optimization level by default. However tests for enqueue_kernel() in clang are compiled with ‘-O0’.</p>
<p class="x_MsoNormal">So, it seems we have a bug here. Am I right?</p>
<p class="x_MsoNormal"> </p>
<p class="x_MsoNormal">Thanks,</p>
<p class="x_MsoNormal">Kristina</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</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.</span></p>
</div>
</div>
</div>
</div>
</div>
</body>
</html>