[Mesa-dev] [Bug 99856] OpenCL Hello world returns "unsupported call to function get_local_size"

[bugzilla-daemon at freedesktop.org]

https://bugs.freedesktop.org/show_bug.cgi?id=99856

--- Comment #14 from Jan Vesely <jv356 at scarletmail.rutgers.edu> ---
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.

-- 
You are receiving this mail because:
You are the QA Contact for the bug.
You are the assignee for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20170225/cc7bf0fa/attachment.html>