diff --git a/libc/src/__support/RPC/CMakeLists.txt b/libc/src/__support/RPC/CMakeLists.txt
--- a/libc/src/__support/RPC/CMakeLists.txt
+++ b/libc/src/__support/RPC/CMakeLists.txt
@@ -2,6 +2,7 @@
   rpc
   HDRS
     rpc.h
+    rpc_util.h
   DEPENDS
     libc.src.__support.common
     libc.src.__support.CPP.atomic
diff --git a/libc/src/__support/RPC/rpc.h b/libc/src/__support/RPC/rpc.h
--- a/libc/src/__support/RPC/rpc.h
+++ b/libc/src/__support/RPC/rpc.h
@@ -18,6 +18,7 @@
 #ifndef LLVM_LIBC_SRC_SUPPORT_RPC_RPC_H
 #define LLVM_LIBC_SRC_SUPPORT_RPC_RPC_H
 
+#include "rpc_util.h"
 #include "src/__support/CPP/atomic.h"
 
 #include <stdint.h>
@@ -101,8 +102,10 @@
   }
   // Wait for the server to work on the buffer and respond.
   if (!in & out) {
-    while (!in)
+    while (!in) {
+      sleep_briefly();
       in = inbox->load(cpp::MemoryOrder::RELAXED);
+    }
     atomic_thread_fence(cpp::MemoryOrder::ACQUIRE);
   }
   // Apply \p use to the buffer and signal the server.
@@ -114,8 +117,10 @@
   }
   // Wait for the server to signal the end of the protocol.
   if (in & !out) {
-    while (in)
+    while (in) {
+      sleep_briefly();
       in = inbox->load(cpp::MemoryOrder::RELAXED);
+    }
     atomic_thread_fence(cpp::MemoryOrder::ACQUIRE);
   }
 }
diff --git a/libc/src/__support/RPC/rpc_util.h b/libc/src/__support/RPC/rpc_util.h
new file mode 100644
--- /dev/null
+++ b/libc/src/__support/RPC/rpc_util.h
@@ -0,0 +1,32 @@
+//===-- Shared memory RPC client / server utilities -------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_SUPPORT_RPC_RPC_UTILS_H
+#define LLVM_LIBC_SRC_SUPPORT_RPC_RPC_UTILS_H
+
+#include "src/__support/macros/attributes.h"
+#include "src/__support/macros/properties/architectures.h"
+
+namespace __llvm_libc {
+namespace rpc {
+
+/// Suspend the thread briefly to assist the thread scheduler during busy loops.
+LIBC_INLINE void sleep_briefly() {
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX) && __CUDA_ARCH__ >= 700
+  asm("nanosleep.u32 64;" ::: "memory");
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+  __builtin_amdgcn_s_sleep(2);
+#else
+  // Simply do nothing if sleeping isn't supported on this platform.
+#endif
+}
+
+} // namespace rpc
+} // namespace __llvm_libc
+
+#endif