(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
signature.asc
Description: PGP signature
_______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/mesa-dev
