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