Hello there,
Has anyone here used the recent LIBCLC (current git version) with the recent git version LLVM (3.9)?
It seems to me that the libraries are not compiled properly and so the generated GPU kernel code does not have the implementation of OpenCl API.
Thanks,
Azin
I built it yesterday. it works. actually, you can only build libclc using llvm 3.9.
yes, the interface of writeBitcodeToFile changed. make sure you delete old ‘*.o’ files before.
“so the generated GPU kernel code does not have the implementation of OpenCl API.”
~~what’s that mean?
It’s not clear what is actually wrong from your original message, I think you need to give some more information as to what you are doing: Example source, what target GPU, compiler error messages or other evidence of “it’s wrong” (llvm IR, disassembly, etc) …
Here is my LLVM version:
LLVM version 3.9.0svn
Optimized build with assertions.
Built Apr 5 2016 (15:33:13).
Default target: x86_64-unknown-linux-gnu
Host CPU: haswell
I am able to ‘configure’ and ‘make’ libclc with llvm-config as well but I am getting a bunch of warnings like following during the make process by llvm-link:
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/integer/add_sat.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/integer/sub_sat.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/subnormal_helper_func.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/atomic/atomic_impl.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/integer/add_sat_impl.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/integer/clz_if.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/integer/clz_impl.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘nvptx–nvidiacl/lib/integer/sub_sat_impl.ll.bc’ is ‘’ whereas ‘llvm-link’ is ‘e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
WARNING: Linking two modules of different data layouts: ‘generic–/lib/subnormal_use_default.bc’ is ‘’ whereas ‘llvm-link’ is 'e-p:32:32-i64:64-v16:16-v32:32-n16:32:64’
Finally I am able to do ‘make install’ without any errors.
I am following steps in here to generate a .bc file for OpenCL kernels; however, after disassembling the linked bc file I see that only the declarations of OpenCl functions has been added and no definition is there.
I suspect that libclc libraries are not compiled completely so that the .bc does not contain any definition from libclc.
Here is .ll version of my linked .bc file (which exactly matches the kernel.ll before linking with libclc):
; ModuleID = ‘kernel.linked.bc’
target datalayout = “e-p:32:32-i64:64-v16:16-v32:32-n16:32:64”
target triple = “nvptx–nvidiacl”
; Function Attrs: noinline nounwind
define void @dataParallel(float addrspace(1)* nocapture readonly %A, float addrspace(1)* nocapture readonly %B, float addrspace(1)* nocapture %C) #0 {
entry:
%call = tail call i32 @get_global_id(i32 0) #2
%mul = shl i32 %call, 2
%arrayidx = getelementptr inbounds float, float addrspace(1)* %A, i32 %mul
%0 = load float, float addrspace(1)* %arrayidx, align 4, !tbaa !9
%arrayidx2 = getelementptr inbounds float, float addrspace(1)* %B, i32 %mul
%1 = load float, float addrspace(1)* %arrayidx2, align 4, !tbaa !9
%add3 = fadd float %0, %1
%arrayidx5 = getelementptr inbounds float, float addrspace(1)* %C, i32 %mul
store float %add3, float addrspace(1)* %arrayidx5, align 4, !tbaa !9
%add6 = or i32 %mul, 1
%arrayidx7 = getelementptr inbounds float, float addrspace(1)* %A, i32 %add6
%2 = load float, float addrspace(1)* %arrayidx7, align 4, !tbaa !9
%arrayidx9 = getelementptr inbounds float, float addrspace(1)* %B, i32 %add6
%3 = load float, float addrspace(1)* %arrayidx9, align 4, !tbaa !9
%sub = fsub float %2, %3
%arrayidx11 = getelementptr inbounds float, float addrspace(1)* %C, i32 %add6
store float %sub, float addrspace(1)* %arrayidx11, align 4, !tbaa !9
%add12 = or i32 %mul, 2
%arrayidx13 = getelementptr inbounds float, float addrspace(1)* %A, i32 %add12
%4 = load float, float addrspace(1)* %arrayidx13, align 4, !tbaa !9
%arrayidx15 = getelementptr inbounds float, float addrspace(1)* %B, i32 %add12
%5 = load float, float addrspace(1)* %arrayidx15, align 4, !tbaa !9
%mul16 = fmul float %4, %5
%arrayidx18 = getelementptr inbounds float, float addrspace(1)* %C, i32 %add12
store float %mul16, float addrspace(1)* %arrayidx18, align 4, !tbaa !9
%add19 = or i32 %mul, 3
%arrayidx20 = getelementptr inbounds float, float addrspace(1)* %A, i32 %add19
%6 = load float, float addrspace(1)* %arrayidx20, align 4, !tbaa !9
%arrayidx22 = getelementptr inbounds float, float addrspace(1)* %B, i32 %add19
%7 = load float, float addrspace(1)* %arrayidx22, align 4, !tbaa !9
%div = fdiv float %6, %7, !fpmath !13
%arrayidx24 = getelementptr inbounds float, float addrspace(1)* %C, i32 %add19
store float %div, float addrspace(1)* %arrayidx24, align 4, !tbaa !9
ret void
}
declare i32 @get_global_id(i32) #1
attributes #0 = { noinline nounwind “disable-tail-calls”=“false” “less-precise-fpmad”=“false” “no-frame-pointer-elim”=“true” “no-frame-pointer-elim-non-leaf” “no-infs-fp-math”=“false” “no-nans-fp-math”=“false” “stack-protector-buffer-size”=“8” “unsafe-fp-math”=“false” “use-soft-float”=“false” }
attributes #1 = { “disable-tail-calls”=“false” “less-precise-fpmad”=“false” “no-frame-pointer-elim”=“true” “no-frame-pointer-elim-non-leaf” “no-infs-fp-math”=“false” “no-nans-fp-math”=“false” “stack-protector-buffer-size”=“8” “unsafe-fp-math”=“false” “use-soft-float”=“false” }
attributes #2 = { nounwind }
Here is my LLVM version:
LLVM version 3.9.0svn
Optimized build with assertions.
Built Apr 5 2016 (15:33:13).
Default target: x86_64-unknown-linux-gnu
Host CPU: haswell
I am able to ‘configure' and ‘make' libclc with llvm-config as well but I am getting a bunch of warnings like following during the make process by llvm-link:
Hi,
What command did you use for linking. In that stack overflow post I
see:
llvm-link libclc/built_libs/nvptx64--nvidiacl.bc test.ll -o test.linked.bc
Which I'm not sure is correct. Try passing test.ll to llvm-link before
the bitcode library.
-Tom