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

[Francisco Jerez]

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?

> 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.

> 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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 227 bytes
Desc: not available
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20180208/7887a213/attachment-0001.sig>