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

[Pierre Moreau]

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

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

Regards,
Pierre
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: not available
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20180208/bec8b1d1/attachment.sig>