[libc] [llvm] [OpenMP] Build OpenMP DeviceRTL on SPIR-V (PR #121600)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload

@llvm/pr-subscribers-libc

Author: Nick Sarnie (sarnex)

<details>
<summary>Changes</summary>

We actually can build it today with the current OMP FE.

We want to minimize device-specific code with SPIR-V, so I used the `spirv64` triple instead of `spirv64-intel` which is currently the only planned SPIR-V based OMP offloading triple. The existing driver code finds the library fine.

I also made a minor change in the RPC code called when building the RTL, `__has_builtin(__builtin_ia32_pause)` was returning true and then we were getting an assert saying it's not supported on this target, it seems there are no features associated with `__builtin_ia32_pause` so `__has_builtin` always returns true on SPIR-V, so just add a arch check so it does nothing as there's no sleep we can use.

---
Full diff: https://github.com/llvm/llvm-project/pull/121600.diff


2 Files Affected:

- (modified) libc/shared/rpc_util.h (+1-1) 
- (modified) offload/DeviceRTL/CMakeLists.txt (+3) 


``````````diff
diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 9406de59f63b71..6b9df43776edac 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -157,7 +157,7 @@ RPC_ATTRS void sleep_briefly() {
     asm("nanosleep.u32 64;" ::: "memory");
 #elif defined(__AMDGPU__) && defined(RPC_TARGET_IS_GPU)
   __builtin_amdgcn_s_sleep(2);
-#elif __has_builtin(__builtin_ia32_pause)
+#elif __has_builtin(__builtin_ia32_pause) && !defined(__SPIRV__)
   __builtin_ia32_pause();
 #elif __has_builtin(__builtin_arm_isb)
   __builtin_arm_isb(0xf);
diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt
index 22940264f9b19a..e707429a5196cf 100644
--- a/offload/DeviceRTL/CMakeLists.txt
+++ b/offload/DeviceRTL/CMakeLists.txt
@@ -264,6 +264,9 @@ compileDeviceRTLLibrary(amdgpu amdgcn-amd-amdhsa -Xclang -mcode-object-version=n
 add_custom_target(omptarget.devicertl.nvptx)
 compileDeviceRTLLibrary(nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx63)
 
+add_custom_target(omptarget.devicertl.spirv64)
+compileDeviceRTLLibrary(spirv64 spirv64)
+
 # Archive all the object files generated above into a static library
 add_library(omptarget.devicertl STATIC)
 set_target_properties(omptarget.devicertl PROPERTIES

``````````

</details>


https://github.com/llvm/llvm-project/pull/121600