From b2d153466d6ea1d29c1075eb27a1fae7ab14bfbe Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Fri, 3 Jan 2025 12:31:00 -0800 Subject: [PATCH 1/3] [OpenMP] Build OpenMP DeviceRTL on SPIRV64 Signed-off-by: Sarnie, Nick --- 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 9406de59f63b7..6b9df43776eda 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 22940264f9b19..e707429a5196c 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" Date: Tue, 7 Jan 2025 14:07:13 -0800 Subject: [PATCH 2/3] Fix __has_builtin workaround Signed-off-by: Sarnie, Nick --- 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 6b9df43776eda..c544037c77578 100644 --- a/libc/shared/rpc_util.h +++ b/libc/shared/rpc_util.h @@ -152,12 +152,14 @@ template 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" Date: Tue, 7 Jan 2025 14:14:15 -0800 Subject: [PATCH 3/3] Update comment Signed-off-by: Sarnie, Nick --- 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 c544037c77578..e264a70b1ef58 100644 --- a/libc/shared/rpc_util.h +++ b/libc/shared/rpc_util.h @@ -153,7 +153,7 @@ template 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");