In a nutshell, this moves our libomptarget code to populate the offload subproject. With this commit, users need to enable the new LLVM/Offload subproject as a runtime in their cmake configuration. No further changes are expected for downstream code. Tests and other components still depend on OpenMP and have also not been renamed. The results below are for a build in which OpenMP and Offload are enabled runtimes. In addition to the pure `git mv`, we needed to adjust some CMake files. Nothing is intended to change semantics. ``` ninja check-offload ``` Works with the X86 and AMDGPU offload tests ``` ninja check-openmp ``` Still works but doesn't build offload tests anymore. ``` ls install/lib ``` Shows all expected libraries, incl. - `libomptarget.devicertl.a` - `libomptarget-nvptx-sm_90.bc` - `libomptarget.rtl.amdgpu.so` -> `libomptarget.rtl.amdgpu.so.18git` - `libomptarget.so` -> `libomptarget.so.18git` Fixes: https://github.com/llvm/llvm-project/issues/75124 --------- Co-authored-by: Saiyedul Islam <Saiyedul.Islam@amd.com>
This patch outlines the SPMD code path into a separate function that can be called directly.
Summary: The LLVM style uses /*Foo=*/ when indicating the name of a constant. See https://llvm.org/docs/CodingStandards.html#comment-formatting. This is useful for consistency, as well as because `clang-format` understands this syntax and formats it more cleanly. Do a bulk update of this syntax.
- [Libomptarget] Make the references to 'malloc' and 'free' weak. - [Libomptarget][NFC] Use C++ style attributes instead
This is good for performance and compile time and the indirection (+ switch statements) is nothing that needs to be preserved.
This feature was supposed to allow you to trace execution inside of Libomptarget. However, this never really worked properly. The printing was always reoganized, only worked for single threads, and pretty much only told you a handful of things about a runtime library that's an implementation detail to all users. Despite this, it contributed about 40% of the total filesize of the deviceRTL. This patch simply removes this functionalit which I think was past due. Reviewed By: tianshilei1992 Differential Revision: https://reviews.llvm.org/D157001
This change makes the naming more consistent, I hope.
If we store a constant in an ICV it is easier for the optimizer to propagate it. Since we often use the full block for the thread limit and the parallel team size, we can instead replace that dynamic value with a constant that otherwise cannot occur, here 0.
We ended up with `llvm.assume(icmp ne ptr as(4) null, as(4) @str)` because the string in address space 4 was not known to be non-null. There is no need to create these assumes.
It turns out that the __builtin_amdgcn_s_barrier() alone does not emit a fence. We somehow got away with this and assumed it would work as it (hopefully) is correct on the NVIDIA path where we just emit a __syncthreads. After talking to @arsenm we now (mostly) align with the OpenCL barrier implementation [1] and emit explicit fences for AMDGPUs. It seems this was the underlying cause for #59759, but I am not 100% certain. There is a chance this simply hides the problem. Fixes: https://github.com/llvm/llvm-project/issues/59759 [1] https://github.com/RadeonOpenCompute/ROCm-Device-Libs/blob/07b347366eb2c6ebc3414af323c623cbbbafc854/opencl/src/workgroup/wgbarrier.cl#L21