I’ve been experimenting with fusing kernels that contain parts written in CUDA/C++ and other parts written in Rust. Unlike clang, rustc does not do early injection of libdevice code so I need to link path/to/libdevice.10.bc manually when invoking llvm-link, and performance is sensitive to LTO. Reading through this thread, it appears there is interest in handling calls to the likes of __nv_log1pf and llvm.sin earlier/more automatically at the IR level (versus clang/lib/Headers/). Has there been upstream work on this (so rustc can/will be able to just use it) or should rustc be handling it manually?
Rust is a good point. Fortran doesn’t get on well with the clang/headers hack either.
Llvm libc has growing GPU support. That might end up the answer to llvm.sin et al, where libc is in bitcode, but doesn’t have it’s own sin implementation yet. It sort of works as another forwarding layer to the vendor lib.
Do you have some thoughts on what a good IR solution would look like from the rust perspective?
I am not convinced the inline functions in a header approach of clang is good or necessary. The right thing seems to be to call sin “sin”, turn it into llvm.sin where flags permit, and lower “sin” and “llvm.sin” in the target back end. Like the non GPU targets do.
1 Like
We have sinf at least thanks to @lntue, but the 64-bit and 16-bit implementations are still in development. The rest right now are a direct re-mapping to __nv_sin if you have it in your system. I’m not fan of this implementation so it’ll be removed once we get more of the core math functions supported.
That route still has many thorns we haven’t traversed yet. One problem is that we still do an eager replacement, so sin will become __nv_sin without -nogpuinc being passed. The second is that llvm.sin intrinsic functions will not be resolved in time for optimal linking (i.e. LTO) or are not currently handled. That quires passing -fno-builitin to prevent the intrinsic conversion. The final problems are caused by the implementations of these functions being freestanding as is discussed in [RFC] libc -ffreestanding / -fno-builtin - #2 by jhuber6. The long story short is that there’s logic that prevents them from being inlined correctly if -fno-builtin isn’t used.
If you work through all of that, the following will work (OpenMP offloading example) and use the implementation of sinf from the LLVM libmgpu.a implementation from the libc for GPUs — The LLVM C Library project.
#include <math.h>
#include <stdio.h>
int main() {
float x = 0.0;
#pragma omp target map(from : x)
{ x = sinf(M_PI / 2.0f); }
printf("%f\n", x);
}
Then compiling and running.
> clang math.c -fopenmp --offload-arch=sm_89 -nogpuinc -fno-builtin -lm -lmgpu
> ./a.out
1.000000
I like the idea of “just” linking math definitions, so hopefully we can iron out these kinks in time.
1 Like
We never upstreamed either solution but IIRC we did implement one of them (@wsmoses can confirm and maybe point you at it).
There are multiple interleaved problems/proposals here and it might be good to clarify which parts you require. Let me try to type some out:
- Handle (math) intrinsics on (GPU) backends that do not support them yet.
- Handle math runtime calls on GPUs.
- Provide portable (math) library functions.
- Replace libdevice with something custom?
There is also the question of your timeline.
Once libm for GPUs is done, which will take a while longer, the driver would link in libm_gpu.a if you (implicitly) specify -lm, either per TU or once late if you do LTO-style linking.
You still want to lower the intrinsics early, probably via [llvm-dev] [RFC] The `implements` attribute, or how to swap functions statically but late.
Together, you can get easy portable math on GPUs but the progress is slow right now.
If you want a solution right now, you can do what we did for Breaking the Vendor Lock - Performance Portable Programming Through OpenMP as Target Independent Runtime Layer (Conference) | OSTI.GOV
Compile the clang/lib/Headers/ into a static library with LLVM-IR targeting all the different GPUs.
You then link against what you need, e.g., with the clang-linker-wrapper, and you can then implement your math functions as if you target libm on the host.
Hope this helps,
~ J
1 Like
Thanks, all. Rust treats nvptx64 as a no_std target so we don’t have access to the standard math operations. (There has been talk of moving it from std to core, but that was years ago and seems stalled indefinitely.) There is a popular crate called num-traits that can transparently bind libm (pure Rust port of MUSL math library, which thus works on all targets) (implementation). I believe the best supported path in the Rust ecosystem will be to add a cuda or device feature to num-traits that would produce symbols like __nv_sinf or target-agnostic names that LLVM can handle.
The question is what to do in rustc to ensure those are resolved. If it’s going to be entirely in the linker, then I think I know what needs to happen, but that relies heavily on LTO for performance. (Notably, the nvptx linker in the Rust repository is broken/disabled, but this PR (embedded-linker) is in review and likely to replace it. It simply calls into LLVM command-line tools at present.)
My current solution just uses a system path for libdevice.10.bc, but this is just for a personal sandox and is not appropriate to get upstream into rustc (which is what I want). I have not yet explored what is needed for amdgpu, though I’ve noticed a similar pattern in clang/lib/Headers/__clang_hip_libdevice_declares.h. I see the code to find a suitable libdevice.bc lives in clang/lib/Driver/ToolChains/{Cuda,AMDGPU}.cpp, though it sounds like nobody here really likes the way Clang handles this so it wouldn’t be a good use of time porting the above to rustc.
If there are no better options, I can wait for the embedded-linker PR above to merge, then port the path search logic from clang and add that to the llvm-link invocation. It looks like the implements attribute hasn’t been implemented and I don’t understand the benefit of compiling the clang/lib/Headers/ into a static library (that would then need to be distributed) versus finding libdevice.10.bc. I may well have missed something so any feedback is welcome and if there can be an upstream LLVM solution with proportional effort, I may be able to contribute.
The former doesn’t eliminate the need for the latter. It simply allows you to lower to target libm functions and get the translation to target specific math calls for AMD and NVIDIA. (The Intel fork might have similar ones for Intel too.) Basically, you’d target sinf as if you had libm, then link the static lib and libdevice, and you’re done. All the headers provide you is the lowering from libm to target_impl.
If you have a portable libm, it might not bee that bad performance-wise, assuming you can link it early or do LTO. See our results for vendor vs LLVM libm’s agnostic impl: https://dl.acm.org/doi/pdf/10.1145/3624062.3624166
I think the opposite of implements was implemented but not upstreamed. Either direction is very straight forward, like a day or so. You’d basically use that to replace the static library discussed above, but that’s it. You still depend on libdevice. The only way to get rid of that is musl/LLVM libm targeting devices directly.
1 Like
Thanks, interesting paper. It may warrant further testing, but rust/musl libm is quite slow compared to standard host implementations so I didn’t expect it to be competitive on device.
In terms of linker requirements, if I write the following in a.cu:
extern "C" __device__ float swirl(float x, float y) {
return log1pf(x) + sinf(y);
}
then
$ clang -S -emit-llvm --cuda-device-only a.cu
produces a a-cuda-*.ll with no dependencies on libdevice, thus JIT also sheds those dependencies so that
$ llvm-link outer_kernel.ll a-cuda-*.ll | opt | llc -o kernel.ptx
is sufficient. This simpler JIT is attractive, especially when running in HPC environments where filesystem metadata is slow/bug-prone. I understand that clang -S still needs libdevice, but that step is AoT for me. (I have two use cases, one doing this JIT fusion of IR from mixed-language sources, the other building kernels AoT in pure Rust.)
This thread started with a proposal to (as I interpret it) remove clang’s eager lowering and have a pass that does the lowering followed immediately by libdevice injection. I assumed locating libdevice would thus move from clang into LLVM (perhaps automatically in a named pass). If that idea is still on the table, then rustc would only need to decide when to run that pass. If that isn’t going to happen, then my understanding is that rustc needs to find libdevice itself and either replicate clang’s injection or add it to the linker. (And finding libdevice would become obsolete once libmgpu is distributed with LLVM, but would still be linked using the usual tools followed by necessary LTO.) Is this accurate?
The motivation here was twofold. For one, we wanted to allow passes to generate math calls and math intrinsic calls late. Second, we wanted to get math optimizations, which trigger for sin and llvm.sin but not for __nv_sin.
If we have a standalone libm, we should be done. If we use libdevice, we need to decide how to link and optimize. That is where the pass comes in as it would allow to get some math optimizations and then replace the calls with target ones that can be inlined.
Thanks, is there a tracking issue to follow libmgpu progress (and does it need some specific help)? In the mean-time, I think I’ll just refine locating libdevice and associated LTO (in my own work as well as trying to get it upstream for Rust).
The main documentation is at libc for GPUs — The LLVM C Library. I haven’t gotten around to writing anything up for the math specific stuff, however the list at Math Functions — The LLVM C Library that @lntue maintains should show all the functions that are implemented in the LLVM libm. However, the GPU implementation chooses to use compiler builtins for a few of them. If you want to help feel free to reach out to me or @lntue.
Ping on this since we have run into the same problem in llvm-flang. Some intrinsics in flang get translated to other dialects which later get translated to LLVM intrinsics. It would be nice to have something working where the library functions and LLVM intrinsics are handled. It seems like this proposal was viewed favorably: NVPTX codegen for llvm.sin (and friends) - #21 by William_Moses3
I would like to get started on this and create PRs unless someone else is already working on this (or some other solution).
Are we talking about making the backend expand things like llvm.sin fully? I know @arsenm has done that for a few math functions in the AMDGPU backend. Similarly there’s been talks about making LLVM intrinsic for every single standard math function.
So, one problem with the intrinsic functions is they’re not implemented the backend will just emit them as a libcall. The GPU builds want to resolve everything via LTO, which is a step too early to resolve llvm.sin with sin or __nv_sin as far as I know.
No, but if the backend does, that should not be disturbed (as it is the “best” solution).
We basically want to have the following flow:
- llvm.math flows in (or is created at some point) which is understood by the optimizer
- libm comes in that implements llvm.math (via attributes)
- backends report what intrinsics they support natively
- a pass in the middle end replaces llvm.math with libm calls
- inliner and IPO have a chance to optimize the call sites
The bespoke metadata, implemented-in-terms-of and so forth, still looks like dead code to me three years after it was the most popular idea.
Write an IR pass that has a table or function names in it. Rows look like llvm.sin, ocml_sin, nvvm_sin.
Use that to replace intrinsics with library calls. Sometimes a vendor specific one, sometimes the libm name. Call it from the backend. Optionally call the inliner afterwards.
Though my recommendation is still to delete the intrinsics. I don’t think the one to one mapping from libm adds any signal.
This is not sufficient due to internalization and the consequent deletion of unused functions. At the time the pass runs, the functions might be gone.
I don’t know what this is supposed to mean. It never made it into upstream, but it is used since the proposal to solve the problem at the time (which still exists).
There is an alternative solution in MLIR for flang, which would be to use the GPU dialect and gpu.module. In mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp there are conversions for some math ops to ocml calls, which would partially solve our problem with intrinsics.
There is also flang/lib/Optimizer/Builder/IntrinsicCall.cpp that lowers intrinsics to either other dialect operations or libm (or other appropriate library) calls. I don’t think these library calls would be converted to ocml library or other vendor specific calls at this point for GPU targets, but could probably be added to the GPUtoROCDL conversion. Perhaps MLIR operations should exist for all library functions so they can be analyzed/optimized in MLIR.
I should add that the GPU module solution is not very modular since knowledge of the math dialect operations are embedded in the gpu dialect (maybe an interface might help). I still think a LLVM intrinsics solution is desirable. One question I have is if two modules have the intrinsic declarations but with different implements metadata, what should happen? It seems like erroring out is reasonable, but is there some existing example where metadata can cause errors?
Easiest is to pick one. Error is fine too. If we had lists, we could keep both.