<html>
<head>
<base href="https://bugs.freedesktop.org/">
</head>
<body>
<p>
<div>
<b><a class="bz_bug_link
bz_status_NEW "
title="NEW - OpenCL Hello world returns "unsupported call to function get_local_size""
href="https://bugs.freedesktop.org/show_bug.cgi?id=99856#c14">Comment # 14</a>
on <a class="bz_bug_link
bz_status_NEW "
title="NEW - OpenCL Hello world returns "unsupported call to function get_local_size""
href="https://bugs.freedesktop.org/show_bug.cgi?id=99856">bug 99856</a>
from <span class="vcard"><a class="email" href="mailto:jv356@scarletmail.rutgers.edu" title="Jan Vesely <jv356@scarletmail.rutgers.edu>"> <span class="fn">Jan Vesely</span></a>
</span></b>
<pre>I think the problem is that the libclc implementation contains uninlined calls
(function calls are not supported an AMD GPUs):
; Function Attrs: alwaysinline nounwind
define linkonce_odr i32 @get_global_id(i32) local_unnamed_addr #12 {
switch i32 %0, label %get_group_id.exit [
i32 0, label %get_group_id.exit.thread
i32 1, label %get_group_id.exit.thread1
i32 2, label %get_group_id.exit.thread2
]
get_group_id.exit.thread: ; preds = %1
%2 = tail call i32 @llvm.amdgcn.workgroup.id.x() #14
%3 = tail call i32 bitcast (i64 (i32)* @get_local_size to i32 (i32)*)(i32 0)
#18
attributes #12 = { alwaysinline nounwind ...
attributes #18 = { nobuiltin nounwind }
libclc commit 520743b generates this code when compiled using llvm-3.9.
This does not happen with LLVM-4/5.</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are the QA Contact for the bug.</li>
<li>You are the assignee for the bug.</li>
</ul>
</body>
</html>