[OpenMP] Build OpenMP DeviceRTL on SPIRV64

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

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);

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