[Nouveau] Some llvm questions (for tgsi backend)
Tom Stellard
tom at stellard.net
Tue Jan 12 07:08:44 PST 2016
On Tue, Jan 12, 2016 at 03:53:10PM +0100, Hans de Goede wrote:
> Hi Tom,
>
> Thanks for taking the time to answer this.
>
> On 11-01-16 18:10, Tom Stellard wrote:
> > On Mon, Jan 11, 2016 at 12:07:14PM +0100, Hans de Goede wrote:
> >> Hi,
> >>
> >> After a few distractions I'm back to work on the llvm tgsi backend. I've
> >> added clang integration and I can now compile a simple opencl program
> >> to something which sort of looks like tgsi.
> >>
> >> You can find my latest work on this here:
> >> http://cgit.freedesktop.org/~jwrdegoede/llvm
> >> http://cgit.freedesktop.org/~jwrdegoede/clang
> >> (the latter may still need to sync)
> >>
> >> I've a little test program of which I have 3 versions now,
> >> 1 raw gallium calls + a tgsi kernel
> >> 2 opencl calls to clover + a tgsi kernel
> >> 3 opencl calls to clover + an opencl kernel
> >>
> >> 1 and 2 have been tested on a kepler card, 3 has been
> >> tested with pocl. My goal for this week is to get
> >> the tgsi backend to produce code which I can copy
> >> and paste into 2 and then have it working on a kepler card.
> >>
> >> The test program looks like this:
> >>
> >> __kernel void test_kern(__global uint *vals, __global uint *buf)
> >> {
> >> uint id = get_global_id(0);
> >>
> >> buf[32 * id] -= vals[id];
> >> }
> >>
> >> The llvm ir looks like this:
> >>
> >> bin/clang -x cl -c -emit-llvm -target tgsi-- -include /usr/share/pocl/include/_kernel.h -o ~/foo.ir -x cl -S ~/foo.cl
> >>
> >> ; ModuleID = '/home/hans/foo.cl'
> >> target datalayout = "E-p:32:32-i64:64:64-f32:32:32-n32"
> >> target triple = "tgsi--"
> >>
> >> ; Function Attrs: nounwind
> >> define void @test_kern(i32 addrspace(1)* nocapture readonly %vals, i32 addrspace(1)* nocapture %buf) #0 {
> >> entry:
> >> %call = tail call i32 @_Z13get_global_idj(i32 0) #2
> >> %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %vals, i32 %call
> >> %0 = load i32, i32 addrspace(1)* %arrayidx, align 4, !tbaa !7
> >> %mul = shl i32 %call, 5
> >> %arrayidx1 = getelementptr inbounds i32, i32 addrspace(1)* %buf, i32 %mul
> >> %1 = load i32, i32 addrspace(1)* %arrayidx1, align 4, !tbaa !7
> >> %sub = sub i32 %1, %0
> >> store i32 %sub, i32 addrspace(1)* %arrayidx1, align 4, !tbaa !7
> >> ret void
> >> }
> >>
> >> declare i32 @_Z13get_global_idj(i32) #1
> >>
> >> attributes #0 = { 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 }
> >>
> >> !opencl.kernels = !{!0}
> >> !llvm.ident = !{!6}
> >>
> >> !0 = !{void (i32 addrspace(1)*, i32 addrspace(1)*)* @test_kern, !1, !2, !3, !4, !5}
> >> !1 = !{!"kernel_arg_addr_space", i32 1, i32 1}
> >> !2 = !{!"kernel_arg_access_qual", !"none", !"none"}
> >> !3 = !{!"kernel_arg_type", !"uint*", !"uint*"}
> >> !4 = !{!"kernel_arg_base_type", !"uint*", !"uint*"}
> >> !5 = !{!"kernel_arg_type_qual", !"", !""}
> >> !6 = !{!"clang version 3.8.0 (http://llvm.org/git/clang.git 9376f992e00569bd08a4ecf3a1d06d8b93c97681) (http://llvm.org/git/llvm.git 7a311143550c6fc01aa5000049825ecc09787440)"}
> >> !7 = !{!8, !8, i64 0}
> >> !8 = !{!"int", !9, i64 0}
> >> !9 = !{!"omnipotent char", !10, i64 0}
> >> !10 = !{!"Simple C/C++ TBAA"}
> >>
> >> And the "tgsi" looks like this:
> >>
> >> .text
> >> .file "/home/hans/foo.cl"
> >> .globl test_kern
> >> test_kern:
> >> BGNSUB
> >> MOVis TEMP1x, 0
> >> CAL _Z13get_global_idj
> >> SHLs TEMP1y, TEMP1x, 7
> >> LOADiis TEMP1z, [4]
> >> UADDs TEMP1y, TEMP1z, TEMP1y
> >> SHLs TEMP1x, TEMP1x, 2
> >> LOADiis TEMP1z, [0]
> >> UADDs TEMP1x, TEMP1z, TEMP1x
> >> LOADgis TEMP1x, [TEMP1x]
> >> INEGs TEMP1x, TEMP1x
> >> LOADgis TEMP1z, [TEMP1y]
> >> UADDs TEMP1x, TEMP1x, TEMP1z
> >> STOREgis [TEMP1y], TEMP1x
> >> RET
> >> ENDSUB
> >>
> >> Working tgsi for this would look like this:
> >>
> >> COMP
> >> DCL SV[0], THREAD_ID[0]
> >> DCL TEMP[0], LOCAL
> >> DCL TEMP[1], LOCAL
> >> IMM UINT32 { 0, 0, 0, 0 }
> >> IMM UINT32 { 4, 0, 0, 0 }
> >> IMM UINT32 { 128, 0, 0, 0 }
> >>
> >> BGNSUB
> >> LOAD TEMP[0].xy, RINPUT, IMM[0]
> >> UMUL TEMP[1].x, SV[0], IMM[1]
> >> UADD TEMP[0].x, TEMP[0], TEMP[1]
> >> UMUL TEMP[1].x, SV[0], IMM[2]
> >> UADD TEMP[0].y, TEMP[0], TEMP[1].xxxx
> >> LOAD TEMP[1].x, RGLOBAL, TEMP[0]
> >> LOAD TEMP[0].x, RGLOBAL, TEMP[0].yyyy
> >> UADD TEMP[1].x, TEMP[0], -TEMP[1]
> >> STORE RGLOBAL.x, TEMP[0].yyyy, TEMP[1]
> >> RET
> >> ENDSUB;
> >>
> >> So my questions (I'm still quite green when it comes to llvm):
> >>
> >> 1) As you can see a proper tgsi program needs a header
> >> to declare which registers (etc) it is using, in which
> >> class-method should I implement this ?
> >>
> >
> > These kinds of things should be emitted by the TGSI implementation of the
> > AsmPrinter class: http://llvm.org/docs/doxygen/html/classllvm_1_1AsmPrinter.html
> > You probably want to use EmitStartOfAsmFile() for the headers.
>
> Ok, is there an easy way to suppress the generation of:
>
> .text
> .file "/home/hans/foo.cl"
> .globl test_kern
>
> ?
>
.file can be avoided pretty easily. You'll have to look through the MC
code to see what bit you need to set to disable it. However, why
do you not want to emit these? I think the easiest think to do is
keep the ELF format and just have clover extract the binary code
from the .text section.
See gallium/drivers/radeon/radeon_elf_util.c
> It seems I need to tackle all 3 of those separately ?
>
> >> 2) Immediates need to be declared with a specific
> >> value and then addressed as IMM[x], how would I go about
> >> this ?
> >>
> >
> > I would recommend adding a pass that replaced immediates in the code with
> > IMM file regitser accesses.
>
> Thinking more about this, I believe I can use the existing const
> mechanism for this (and then add a const generation pass to AsmPrinter).
>
I would recommend not doing this kind of thing in the AsmPrinter if you
can avoid it. I really think you should be able to do this in a
MachineInstr pass.
> How do I tell llvm to never generate instructions using immediates
> and to instead always use the const address space for this ?
>
LLVM usually won't put immediates into instructions unless you have a
TableGen pattern that tells it to do this. I suspect you probably have
some patterns like this.
> >> 3) The get_global_id call needs to be translated into
> >> simply using the SV[0] "register", how would I go about
> >> this ?
> >>
> >
> > get_global_id should be implemented in libclc with tgsi specific intrinsics
> > that read from the system value registers.
>
> I see I will investigate this further.
>
> >> 4) The global and input load / stores are not handled
> >> correctly, I see that the LOAD instructions get postfixed
> >> with a i reps. g for input / global how would I go about
> >> modifying the code emitter (AsmPrinter?) to change "LOADi"
> >> into "LOAD <dest> RINPUT <offset>"?
> >>
> >
> > I don't understand exactly why you want to change this. Are you trying to
> > make the assembly string for input loads and global loads look the same?
>
> Currently the following code is generated:
>
> LOADiis TEMP[1].z, [0]
>
> This should be:
>
> LOAD TEMP[1].z, RINPUT, 0
>
> So I need to add handling to generate the RINPUT here. I'm wondering
> what the best place is for this. Do I simply intercept the LOAD in
> LowerMachineInstrToMCInst and add an extra operand there ?
>
It seems like you just need to change your instruction definitions an/or
assembly printing here.
>
> >> 5) Talking about the lowecase suffixes to the instructions,
> >> these should not be part of the output, how do I filter these?
> >>
> >> 6) And finally, the current llvm-tgsi output uses e.g.
> >> TEMP1y where as for the destination it should use TEMP[1].y
> >> and for the sources it should use TEMP[1].xxxx (so include
> >> proper swizzling info).
> >>
> >
> > You just need to change how the registers are printed in the
> > TGSIRegisterInfo.td file which should fix this.
>
> OK, so I've managed (easy) to change things so that now I get
> (i.e.) :
>
> UADDs TEMP[1].x, TEMP[1].x, TEMP[1].z
>
> This is however not correct TGSI, correct would be:
>
Are you sure the above isn't correct?
> UADD TEMP[1].x, TEMP[1].xxxx, TEMP[1].zzzz
>
> Which may be shortend to:
>
> UADD TEMP[1].x, TEMP[1], TEMP[1].zzzz
>
> TGSI will automatically use the dest vector component from
> the sources, and if you want to use another component you
> need to specify a swizzle order, which is the 4 components
> in an alternate order.
>
> I'm not quite sure yet how to deal with this.
>
These kinds of things can be complicated, I think this is a
good reason to start emitting binary code from LLVM rather than
text. You can always clean up the assembly format later.
-Tom
> Regards,
>
> Hans
More information about the Nouveau
mailing list