[libc-commits] [PATCH] D147238: [libc] Support suspending threads during RPC spin loops

[Joseph Huber via Phabricator via libc-commits]

jhuber6 created this revision.
jhuber6 added reviewers: JonChesterfield, tra, jdoerfert, tianshilei1992, sivachandra.
Herald added subscribers: libc-commits, ecnelises, tschuett.
Herald added projects: libc-project, All.
jhuber6 requested review of this revision.

The RPC interface relies on waiting on atomic signals to coordiante
which side of the protocol is in control of the shared buffer. This
requires waiting on an atomic value until it changes. The GPU client
supports briefly suspending the executing thread group. This is
important for multi-threaded implementations as it allows. The thread
scheduler in the CUs are not fair and may repeatedly choose to run a
thread group that cannot make progress. Sleeping signals that the
scheduler should pick other threads to run so we can guaruntee forward
progress.

This is currently only relevant on the client-side. We could use an
alternative implementation on the server that uses the standard
`nanosleep` on supported hosts.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D147238

Files:
  libc/src/__support/RPC/CMakeLists.txt
  libc/src/__support/RPC/rpc.h
  libc/src/__support/RPC/rpc_util.h


Index: libc/src/__support/RPC/rpc_util.h
===================================================================
--- /dev/null
+++ 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
Index: libc/src/__support/RPC/rpc.h
===================================================================
--- libc/src/__support/RPC/rpc.h
+++ 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) {
       in = inbox->load(cpp::MemoryOrder::RELAXED);
+      sleep_briefly();
+    }
     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) {
       in = inbox->load(cpp::MemoryOrder::RELAXED);
+      sleep_briefly();
+    }
     atomic_thread_fence(cpp::MemoryOrder::ACQUIRE);
   }
 }
Index: libc/src/__support/RPC/CMakeLists.txt
===================================================================
--- libc/src/__support/RPC/CMakeLists.txt
+++ 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


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D147238.509683.patch
Type: text/x-patch
Size: 2565 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/libc-commits/attachments/20230330/d9394601/attachment.bin>