diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 687814b7ff2ae..1ad610662e0c1 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -152,7 +152,9 @@ template <typename T> class optional {
 
 /// Suspend the thread briefly to assist the thread scheduler during busy loops.
 RPC_ATTRS void sleep_briefly() {
-#if __has_builtin(__nvvm_reflect)
+#if defined(__SPIRV__)
+  // FIXME: __has_builtin does not work on offloading targets.
+#elif __has_builtin(__nvvm_reflect)
   if (__nvvm_reflect("__CUDA_ARCH") >= 700)
     asm("nanosleep.u32 64;" ::: "memory");
 #elif __has_builtin(__builtin_amdgcn_s_sleep)
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