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

[Jan Vesely]

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. SPIR-V binary can also call exp2(fp64), which does
not have an equivalent GPU instruction.
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).

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. 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: 488 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20180208/46ca2707/attachment-0001.sig>