https://bugs.freedesktop.org/show_bug.cgi?id=99856
--- Comment #14 from Jan Vesely <jv...@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._______________________________________________
mesa-dev mailing list
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
