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

[Nick Sarnie via libc-commits]

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

>From b2d153466d6ea1d29c1075eb27a1fae7ab14bfbe Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sarnie at intel.com>
Date: Fri, 3 Jan 2025 12:31:00 -0800
Subject: [PATCH 1/3] [OpenMP] Build OpenMP DeviceRTL on SPIRV64

Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>
---
 libc/shared/rpc_util.h           | 2 +-
 offload/DeviceRTL/CMakeLists.txt | 3 +++
 2 files changed, 4 insertions(+), 1 deletion(-)

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

>From c40b5b625142798acf94de93645ca1df143798d7 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sarnie at intel.com>
Date: Tue, 7 Jan 2025 14:07:13 -0800
Subject: [PATCH 2/3] Fix __has_builtin workaround

Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>
---
 libc/shared/rpc_util.h | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 6b9df43776edac..c544037c775789 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -152,12 +152,14 @@ template <typename T> class optional {
 
 /// Suspend the thread briefly to assist the thread scheduler during busy loops.
 RPC_ATTRS void sleep_briefly() {
-#if defined(__NVPTX__) && defined(RPC_TARGET_IS_GPU)
+#if defined(__SPIRV__)
+  // It's unsupported and __has_builtin doesn't always work as we expect.
+#elif defined(__NVPTX__) && defined(RPC_TARGET_IS_GPU)
   if (__nvvm_reflect("__CUDA_ARCH") >= 700)
     asm("nanosleep.u32 64;" ::: "memory");
 #elif defined(__AMDGPU__) && defined(RPC_TARGET_IS_GPU)
   __builtin_amdgcn_s_sleep(2);
-#elif __has_builtin(__builtin_ia32_pause) && !defined(__SPIRV__)
+#elif __has_builtin(__builtin_ia32_pause)
   __builtin_ia32_pause();
 #elif __has_builtin(__builtin_arm_isb)
   __builtin_arm_isb(0xf);

>From f0cc084c89d5098e2127b5f36dbffec556aaaf04 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sarnie at intel.com>
Date: Tue, 7 Jan 2025 14:14:15 -0800
Subject: [PATCH 3/3] Update comment

Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>
---
 libc/shared/rpc_util.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index c544037c775789..e264a70b1ef58c 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -153,7 +153,7 @@ template <typename T> class optional {
 /// Suspend the thread briefly to assist the thread scheduler during busy loops.
 RPC_ATTRS void sleep_briefly() {
 #if defined(__SPIRV__)
-  // It's unsupported and __has_builtin doesn't always work as we expect.
+  // FIXME: __has_builtin does not work on offloading targets.
 #elif defined(__NVPTX__) && defined(RPC_TARGET_IS_GPU)
   if (__nvvm_reflect("__CUDA_ARCH") >= 700)
     asm("nanosleep.u32 64;" ::: "memory");