[Mesa-dev] Adding a SPIR-V target to libclc

Jan Vesely jan.vesely at rutgers.edu
Mon Feb 19 21:33:11 UTC 2018 

On Thu, 2018-02-15 at 22:50 -0800, Francisco Jerez wrote:
> Jan Vesely <jan.vesely at rutgers.edu> writes:
> 
> > On Thu, 2018-02-15 at 20:36 -0800, Francisco Jerez wrote:
> > > Jan Vesely <jan.vesely at rutgers.edu> writes:
> > > 
> > > > On Thu, 2018-02-08 at 15:56 -0800, Francisco Jerez wrote:
> > > > > Jan Vesely <jan.vesely at rutgers.edu> writes:
> > > > > 
> > > > > > On Thu, 2018-02-08 at 23:16 +0100, Pierre Moreau wrote:
> > > > > > > (Moving the conversation to its own thread.)
> > > > > > > 
> > > > > > > > target agnostic libclc is rather difficult to do. CLC includes 3 levels
> > > > > > > > of precision on float (fp32) operands; full, half, native. The
> > > > > > > > implementation of each depends on capabilites of specific device (e.g.
> > > > > > > > vega(VI+?) can do 1 ULP log2/exp2 in hw, other targets need sw
> > > > > > > > implementation to meet CLC requirement of 3ulp). Any conversion backend
> > > > > > > > would thus need to implement sw versions of math builtins for targets
> > > > > > > > that can't perform the op in HW.
> > > > > > > 
> > > > > > > My initial thought for the target agnostic libclc, was to just provide some
> > > > > > > (fake?) implementations of OpenCL built-in functions to make clang happy and
> > > > > > > let me compile kernels using “get_global_id()”, as well as include headers
> > > > > > > defining OpenCL specific types like “float4” or others. If there is another
> > > > > > > (better?) way to achieve this, I am all ears. (There is probably one, as I had
> > > > > > > no issues when using the Khronos LLVM/clang fork rather than Tomeu’s
> > > > > > > out-of-tree module, the former having also some bits and pieces in clang.)
> > > > > > 
> > > > > > I don't think you need libclc for this. workitem IDs are
> > > > > > platform/device specific, and iiuc SPIR-V builtins should handle it in
> > > > > > an abstract way [0]. any conversion consuming SPIR-V needs to replace
> > > > > > those with device/platform specific way of obtaining the information.
> > > > > > you can also use clang's clc header to provide data types [1].
> > > > > > 
> > > > > > 
> > > > > > [0] https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#B
> > > > > > uiltIn
> > > > > > [1] https://github.com/llvm-mirror/clang/blob/master/lib/Headers/opencl
> > > > > > -c.h
> > > > > > 
> > > > > > > 
> > > > > > > > Extending the current libclc to provide target specific SPIR-V binaries
> > > > > > > > in addition to/in place of LLVM IR is rather straightforward. Adding
> > > > > > > > additional targets it's more work since it relies on clang to support
> > > > > > > > those targets.
> > > > > > > 
> > > > > > > I’m curious how those target specific SPIR-V binaries would look like. I can
> > > > > > > imagine how some functions like “OpSign” could be implemented using other
> > > > > > > SPIR-V functions, but how would you handle something like “get_local_id()”? If
> > > > > > > you define it as the built-in “LocalInvocationId” and don’t supply an
> > > > > > > implementation of it, then you lose the target specificness. On the other hand,
> > > > > > > if you want to keep it device-specific, how would you express that in SPIR-V?
> > > > > > 
> > > > > > getting IDs is not a problem. SPIR-V should provide builtins for that.
> > > > > > 
> > > > > > The problem I had in mind is when SPIR-V binary calls e.g. exp2(). You
> > > > > > can either assume that the op needs CLC precision (3 ulp), or device
> > > > > > native precision.
> > > > > 
> > > > > That's up to the SPIR-V extended instruction set specification to define
> > > > > what precision the exp2 built-in is supposed to have.  
> > > > > 
> > > > > > SPIR-V binary can also call exp2(fp64), which does not have an
> > > > > > equivalent GPU instruction.
> > > > > 
> > > > > Then it should probably be lowered by the SPIR-V front-end, right?
> > > > 
> > > > I'm not sure what you mean by "spir-v frontend". If it's the tool that
> > > > generates SPIR-V, then no, not really. 
> > > 
> > > No, I meant the SPIR-V front-end of the driver (or whatever translation
> > > pass in control of the driver is translating machine-agnostic SPIR-V
> > > into some other more hardware-specific representation of the program).
> > 
> > OK. my question still stands. How does generic SPIR-V based libclc
> > help the process?
> > 
> 
> That I can think of now, it would remove the need for maintaining any
> target-specific knowledge in libclc, for plumbing target-specific
> information in order to select the right libclc flavour at link time,

it would only move the specific decisions to SPIR-V lowering time. I
understand the advantage of cross language usefulness, but I'm not sure
how practical it is.
Taking the below example of exp2(fp64). CLC requires precision <= 2ulp,
other languages might have different requirements. Thus to achieve good
performance, you'd need to lower to different routine for each
precision requirement.

> and it would allow solving common problems in a place where there is a
> chance that the solution could be shared among different drivers and
> APIs (e.g. the exp2(fp64) lowering example you mentioned earlier is not
> exclusively useful to CL).

the exp2 example is not something that could be addressed in generic
SPIR-V libclc, since the decision is hw specific. Sure we can provide
implementation of all CLC builtins using only the core SPIR-V
operations, but if a SPIR-V input uses clc extended instructions the
same functionality would have to be implemented in SPIR-V lowering
anyway, so it's just simpler to implement libclc as single op wrappers
over CLC extended ops. Am I missing anything?

Jan

> 
> > Jan
> > 
> > > 
> > > > My understanding is that those are run prior to application
> > > > distribution, and therefore have no information about the target HW.
> > > > 
> > > > So if a program imports "CLC.std.11" extended instruction set to get
> > > > access CLC builtin functions. What would a generic SPIR-V libclc
> > > > provide?
> > > > 
> > > > > 
> > > > > > It's easier to translate these to libclc function calls (combined with
> > > > > > the right library implementation of the exp2 builtin), than try to
> > > > > > generate exp2 algorithm when converting to NIR (or anything else
> > > > > > really).
> > > > > > 
> > > > > 
> > > > > But the SPIR-V front-end will need to lower that in terms of
> > > > > instructions supported by the back-end anyway in order to be able to
> > > > > handle general SPIR-V shaders as input, right?  So why re-implement the
> > > > > lowering for those operations in libclc in a way that's only going to be
> > > > > useful for the OpenCL C language but not for other APIs?
> > > > > 
> > > > > > The current libclc mostly assumes that LLVM ops are done in device
> > > > > > native precision, and provides sw implementation of operations that
> > > > > > don't have conformant device instruction.
> > > > > 
> > > > > But I don't think there is any disadvantage from having a libclc
> > > > > implementation that doesn't make any precision assumptions beyond what
> > > > > is stated in the SPIR-V spec.  In fact that would have the IMO more
> > > > > desirable advantage that you could re-use one and the same libclc
> > > > > implementation for *all* back-ends that want SPIR-V as input.
> > > > 
> > > > Sure, a compiler-rt library would be more useful (usable by multiple
> > > > languages). However, unlike target specific libclc, it's not available
> > > > atm.
> > > > 
> > > > Jan
> > > > 
> > > > > 
> > > > > > This obviates the need for compiler-rt library.  And alternative
> > > > > > approach is to assume that the ops provide full precision and use
> > > > > > target intrinsics for native precision. it's still target specific if
> > > > > > a library call uses the former or the latter.
> > > > > > 
> > > > > > regards,
> > > > > > Jan
> > > > > > 
> > > > > > > 
> > > > > > > Regards,
> > > > > > > Pierre
> > > > > 
> > > > > _______________________________________________
> > > > > mesa-dev mailing list
> > > > > mesa-dev at lists.freedesktop.org
> > > > > https://lists.freedesktop.org/mailman/listinfo/mesa-dev
> > > > 
> > > > -- 
> > > > Jan Vesely <jan.vesely at rutgers.edu>
> > 
> > -- 
> > Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 488 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20180219/189718d0/attachment.sig>

More information about the mesa-dev mailing list