diff --git a/libc/docs/gpu/rpc.rst b/libc/docs/gpu/rpc.rst
index c87d60640794..bd82e0942f9c 100644
--- a/libc/docs/gpu/rpc.rst
+++ b/libc/docs/gpu/rpc.rst
@@ -117,7 +117,6 @@ done. It can be omitted if asynchronous execution is desired.
       buffer->data[0] = reinterpret_cast<uintptr_t>(fn);
     });
     port.recv([](rpc::Buffer *, uint32_t) {});
-    port.close();
   }
 
 Server Example
@@ -162,7 +161,6 @@ data.
       port->recv([](rpc::Buffer *) {});
       break;
     }
-    port->close();
   }
 
 Function Dispatch
@@ -199,7 +197,6 @@ than submitting asynchronously.
       port->recv([](rpc::Buffer *) {});
       break;
     }
-    port->close();
   }
 
 
@@ -275,7 +272,6 @@ but the following example shows how it can be used by a standard user.
 
       // Only available in-tree from the 'libc' sources.
       handle_libc_opcodes(*port, warp_size);
-      port->close();
     } while (cudaStreamQuery(stream) == cudaErrorNotReady);
   }
 
diff --git a/libc/shared/rpc.h b/libc/shared/rpc.h
index 1233f4c0bbdf..9465868f21fc 100644
--- a/libc/shared/rpc.h
+++ b/libc/shared/rpc.h
@@ -297,7 +297,7 @@ template <bool T> struct Port {
                  uint32_t index, uint32_t out)
       : process(process), lane_mask(lane_mask), lane_size(lane_size),
         index(index), out(out), receive(false), owns_buffer(true) {}
-  RPC_ATTRS ~Port() = default;
+  RPC_ATTRS ~Port() { close(); }
 
 private:
   RPC_ATTRS Port(const Port &) = delete;
@@ -332,6 +332,7 @@ public:
     return lane_mask;
   }
 
+private:
   RPC_ATTRS void close() {
     // Wait for all lanes to finish using the port.
     rpc::sync_lane(lane_mask);
@@ -343,7 +344,6 @@ public:
     process.unlock(lane_mask, index);
   }
 
-private:
   Process<T> &process;
   uint64_t lane_mask;
   uint32_t lane_size;
diff --git a/libc/shared/rpc_dispatch.h b/libc/shared/rpc_dispatch.h
index 1a385c1b7d82..ff4f357378f2 100644
--- a/libc/shared/rpc_dispatch.h
+++ b/libc/shared/rpc_dispatch.h
@@ -208,7 +208,6 @@ dispatch(rpc::Client &client, FnTy, CallArgs... args) {
   using BufferTy = rpc::conditional_t<rpc::is_void_v<RetTy>, uint8_t, RetTy>;
   BufferTy ret{};
   port.recv_n(&ret);
-  port.close();
 
   if constexpr (!rpc::is_void_v<RetTy>)
     return ret;
diff --git a/libc/src/__support/GPU/allocator.cpp b/libc/src/__support/GPU/allocator.cpp
index f8a3b464ea22..588cf3277646 100644
--- a/libc/src/__support/GPU/allocator.cpp
+++ b/libc/src/__support/GPU/allocator.cpp
@@ -61,7 +61,6 @@ static void *rpc_allocate(uint64_t size) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ptr = reinterpret_cast<void *>(buffer->data[0]);
       });
-  port.close();
   return ptr;
 }
 
@@ -71,7 +70,6 @@ static void rpc_free(void *ptr) {
   port.send([=](rpc::Buffer *buffer, uint32_t) {
     buffer->data[0] = reinterpret_cast<uintptr_t>(ptr);
   });
-  port.close();
 }
 
 // Convert a potentially disjoint bitmask into an increasing integer per-lane
diff --git a/libc/src/__support/OSUtil/gpu/exit.cpp b/libc/src/__support/OSUtil/gpu/exit.cpp
index 85f8183aafa9..1024fcb8276b 100644
--- a/libc/src/__support/OSUtil/gpu/exit.cpp
+++ b/libc/src/__support/OSUtil/gpu/exit.cpp
@@ -24,7 +24,6 @@ namespace internal {
   port.send([&](rpc::Buffer *buffer, uint32_t) {
     reinterpret_cast<uint32_t *>(buffer->data)[0] = status;
   });
-  port.close();
 
   gpu::end_program();
 }
diff --git a/libc/src/__support/OSUtil/gpu/io.cpp b/libc/src/__support/OSUtil/gpu/io.cpp
index 5d107ab50ae8..34cd4f836204 100644
--- a/libc/src/__support/OSUtil/gpu/io.cpp
+++ b/libc/src/__support/OSUtil/gpu/io.cpp
@@ -18,7 +18,6 @@ void write_to_stderr(cpp::string_view msg) {
   rpc::Client::Port port = rpc::client.open<LIBC_WRITE_TO_STDERR>();
   port.send_n(msg.data(), msg.size());
   port.recv([](rpc::Buffer *, uint32_t) { /* void */ });
-  port.close();
 }
 
 } // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/stdio/gpu/clearerr.cpp b/libc/src/stdio/gpu/clearerr.cpp
index 5a0ca52e33fa..200fc5c2e965 100644
--- a/libc/src/stdio/gpu/clearerr.cpp
+++ b/libc/src/stdio/gpu/clearerr.cpp
@@ -21,7 +21,6 @@ LLVM_LIBC_FUNCTION(void, clearerr, (::FILE * stream)) {
         buffer->data[0] = file::from_stream(stream);
       },
       [&](rpc::Buffer *, uint32_t) {});
-  port.close();
 }
 
 } // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/stdio/gpu/fclose.cpp b/libc/src/stdio/gpu/fclose.cpp
index 1e00515e0650..a2458273bed6 100644
--- a/libc/src/stdio/gpu/fclose.cpp
+++ b/libc/src/stdio/gpu/fclose.cpp
@@ -22,7 +22,6 @@ LLVM_LIBC_FUNCTION(int, fclose, (::FILE * stream)) {
   port.send_and_recv(
       [=](rpc::Buffer *buffer, uint32_t) { buffer->data[0] = file; },
       [&](rpc::Buffer *buffer, uint32_t) { ret = buffer->data[0]; });
-  port.close();
 
   if (ret != 0)
     return EOF;
diff --git a/libc/src/stdio/gpu/feof.cpp b/libc/src/stdio/gpu/feof.cpp
index 3ae308bad60b..20e8583e4c5b 100644
--- a/libc/src/stdio/gpu/feof.cpp
+++ b/libc/src/stdio/gpu/feof.cpp
@@ -24,7 +24,6 @@ LLVM_LIBC_FUNCTION(int, feof, (::FILE * stream)) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ret = static_cast<int>(buffer->data[0]);
       });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/ferror.cpp b/libc/src/stdio/gpu/ferror.cpp
index 64d62e706525..628f37eb77e1 100644
--- a/libc/src/stdio/gpu/ferror.cpp
+++ b/libc/src/stdio/gpu/ferror.cpp
@@ -24,7 +24,6 @@ LLVM_LIBC_FUNCTION(int, ferror, (::FILE * stream)) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ret = static_cast<int>(buffer->data[0]);
       });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/fflush.cpp b/libc/src/stdio/gpu/fflush.cpp
index 0b6ef92f346a..3032554cd11c 100644
--- a/libc/src/stdio/gpu/fflush.cpp
+++ b/libc/src/stdio/gpu/fflush.cpp
@@ -24,7 +24,6 @@ LLVM_LIBC_FUNCTION(int, fflush, (::FILE * stream)) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ret = static_cast<int>(buffer->data[0]);
       });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/fgets.cpp b/libc/src/stdio/gpu/fgets.cpp
index e1c6088b6008..8b2c71ebd106 100644
--- a/libc/src/stdio/gpu/fgets.cpp
+++ b/libc/src/stdio/gpu/fgets.cpp
@@ -31,7 +31,6 @@ LLVM_LIBC_FUNCTION(char *, fgets,
   });
   port.recv_n(&buf, &recv_size,
               [&](uint64_t) { return reinterpret_cast<void *>(str); });
-  port.close();
 
   if (recv_size == 0)
     return nullptr;
diff --git a/libc/src/stdio/gpu/file.h b/libc/src/stdio/gpu/file.h
index 102c6c601b30..ca1bb3b49148 100644
--- a/libc/src/stdio/gpu/file.h
+++ b/libc/src/stdio/gpu/file.h
@@ -63,7 +63,6 @@ LIBC_INLINE uint64_t write_impl(::FILE *file, const void *data, size_t size) {
   port.recv([&](rpc::Buffer *buffer, uint32_t) {
     ret = reinterpret_cast<uint64_t *>(buffer->data)[0];
   });
-  port.close();
   return ret;
 }
 
@@ -86,7 +85,6 @@ LIBC_INLINE uint64_t read_from_stream(::FILE *file, void *buf, size_t size) {
   });
   port.recv_n(&buf, &recv_size, [&](uint64_t) { return buf; });
   port.recv([&](rpc::Buffer *buffer, uint32_t) { ret = buffer->data[0]; });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/fopen.cpp b/libc/src/stdio/gpu/fopen.cpp
index eee3edaac1c1..cb1bff267f80 100644
--- a/libc/src/stdio/gpu/fopen.cpp
+++ b/libc/src/stdio/gpu/fopen.cpp
@@ -26,7 +26,6 @@ LLVM_LIBC_FUNCTION(::FILE *, fopen,
         inline_memcpy(buffer->data, mode, internal::string_length(mode) + 1);
       },
       [&](rpc::Buffer *buffer, uint32_t) { file = buffer->data[0]; });
-  port.close();
 
   return reinterpret_cast<FILE *>(file);
 }
diff --git a/libc/src/stdio/gpu/fseek.cpp b/libc/src/stdio/gpu/fseek.cpp
index 468d65ff77e2..6d24e98167d3 100644
--- a/libc/src/stdio/gpu/fseek.cpp
+++ b/libc/src/stdio/gpu/fseek.cpp
@@ -26,7 +26,6 @@ LLVM_LIBC_FUNCTION(int, fseek, (::FILE * stream, long offset, int whence)) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ret = static_cast<int>(buffer->data[0]);
       });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/ftell.cpp b/libc/src/stdio/gpu/ftell.cpp
index 7ee33d7b05d9..f43c3d8a9b40 100644
--- a/libc/src/stdio/gpu/ftell.cpp
+++ b/libc/src/stdio/gpu/ftell.cpp
@@ -24,7 +24,6 @@ LLVM_LIBC_FUNCTION(long, ftell, (::FILE * stream)) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ret = static_cast<long>(buffer->data[0]);
       });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/remove.cpp b/libc/src/stdio/gpu/remove.cpp
index 3cd72745cb6d..52e15629394d 100644
--- a/libc/src/stdio/gpu/remove.cpp
+++ b/libc/src/stdio/gpu/remove.cpp
@@ -22,7 +22,6 @@ LLVM_LIBC_FUNCTION(int, remove, (const char *path)) {
   port.recv([&](rpc::Buffer *buffer, uint32_t) {
     ret = static_cast<int>(buffer->data[0]);
   });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/rename.cpp b/libc/src/stdio/gpu/rename.cpp
index d7c71fc82682..e4d1048b6fbe 100644
--- a/libc/src/stdio/gpu/rename.cpp
+++ b/libc/src/stdio/gpu/rename.cpp
@@ -23,7 +23,6 @@ LLVM_LIBC_FUNCTION(int, rename, (const char *oldpath, const char *newpath)) {
   port.recv([&](rpc::Buffer *buffer, uint32_t) {
     ret = static_cast<int>(buffer->data[0]);
   });
-  port.close();
 
   return ret;
 }
diff --git a/libc/src/stdio/gpu/ungetc.cpp b/libc/src/stdio/gpu/ungetc.cpp
index fadd1d7db13c..f9ac50e02e2b 100644
--- a/libc/src/stdio/gpu/ungetc.cpp
+++ b/libc/src/stdio/gpu/ungetc.cpp
@@ -25,7 +25,6 @@ LLVM_LIBC_FUNCTION(int, ungetc, (int c, ::FILE *stream)) {
       [&](rpc::Buffer *buffer, uint32_t) {
         ret = static_cast<int>(buffer->data[0]);
       });
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdio/gpu/vfprintf_utils.h b/libc/src/stdio/gpu/vfprintf_utils.h
index 6df46489d6c9..70eb98c838be 100644
--- a/libc/src/stdio/gpu/vfprintf_utils.h
+++ b/libc/src/stdio/gpu/vfprintf_utils.h
@@ -51,7 +51,6 @@ LIBC_INLINE int vfprintf_impl(::FILE *__restrict file,
     port.send_n(str, size);
   }
 
-  port.close();
   return ret;
 }
 
diff --git a/libc/src/stdlib/gpu/abort.cpp b/libc/src/stdlib/gpu/abort.cpp
index 8a7e783990d6..05bd13f3b979 100644
--- a/libc/src/stdlib/gpu/abort.cpp
+++ b/libc/src/stdlib/gpu/abort.cpp
@@ -21,7 +21,6 @@ LLVM_LIBC_FUNCTION(void, abort, ()) {
   port.send_and_recv([](rpc::Buffer *, uint32_t) {},
                      [](rpc::Buffer *, uint32_t) {});
   port.send([&](rpc::Buffer *, uint32_t) {});
-  port.close();
 
   gpu::end_program();
 }
diff --git a/libc/src/stdlib/gpu/system.cpp b/libc/src/stdlib/gpu/system.cpp
index 1677e6051c5f..63f48aa0bf07 100644
--- a/libc/src/stdlib/gpu/system.cpp
+++ b/libc/src/stdlib/gpu/system.cpp
@@ -22,7 +22,6 @@ LLVM_LIBC_FUNCTION(int, system, (const char *command)) {
   port.recv([&](rpc::Buffer *buffer, uint32_t) {
     ret = static_cast<int>(buffer->data[0]);
   });
-  port.close();
 
   return ret;
 }
diff --git a/libc/test/integration/startup/gpu/rpc_interface_test.cpp b/libc/test/integration/startup/gpu/rpc_interface_test.cpp
index eed9b67ef8db..0b7ff549fe39 100644
--- a/libc/test/integration/startup/gpu/rpc_interface_test.cpp
+++ b/libc/test/integration/startup/gpu/rpc_interface_test.cpp
@@ -53,7 +53,6 @@ static void test_interface(bool end_with_send) {
     port.recv([&](LIBC_NAMESPACE::rpc::Buffer *buffer, uint32_t) {
       cnt = buffer->data[0];
     });
-  port.close();
 
   ASSERT_TRUE(cnt == 9 && "Invalid number of increments");
 }
diff --git a/libc/test/integration/startup/gpu/rpc_lane_test.cpp b/libc/test/integration/startup/gpu/rpc_lane_test.cpp
index e80b438829b4..486f9de3c9ff 100644
--- a/libc/test/integration/startup/gpu/rpc_lane_test.cpp
+++ b/libc/test/integration/startup/gpu/rpc_lane_test.cpp
@@ -23,7 +23,6 @@ static void test_add() {
       [&](LIBC_NAMESPACE::rpc::Buffer *buffer, uint32_t) {
         cnt = reinterpret_cast<uint64_t *>(buffer->data)[0];
       });
-  port.close();
   EXPECT_EQ(cnt, gpu::get_lane_id() + 1);
   EXPECT_EQ(gpu::get_thread_id(), gpu::get_lane_id());
 }
diff --git a/libc/test/integration/startup/gpu/rpc_stream_test.cpp b/libc/test/integration/startup/gpu/rpc_stream_test.cpp
index b8c37926d2d4..bb7b1eb3a3db 100644
--- a/libc/test/integration/startup/gpu/rpc_stream_test.cpp
+++ b/libc/test/integration/startup/gpu/rpc_stream_test.cpp
@@ -39,7 +39,6 @@ static void test_stream() {
   port.send_n(send_ptr, send_size);
   port.recv_n(&recv_ptr, &recv_size,
               [](uint64_t size) { return malloc(size); });
-  port.close();
   ASSERT_TRUE(inline_memcmp(recv_ptr, str, recv_size) == 0 && "Data mismatch");
   ASSERT_TRUE(recv_size == send_size && "Data size mismatch");
 
@@ -83,7 +82,6 @@ static void test_divergent() {
   port.send_n(buffer, offset);
   inline_memset(buffer, 0, offset);
   port.recv_n(&recv_ptr, &recv_size, [&](uint64_t) { return buffer; });
-  port.close();
 
   ASSERT_TRUE(inline_memcmp(recv_ptr, &data[offset], recv_size) == 0 &&
               "Data mismatch");
diff --git a/libc/test/integration/startup/gpu/rpc_test.cpp b/libc/test/integration/startup/gpu/rpc_test.cpp
index d46a1adf2857..018745a9e29c 100644
--- a/libc/test/integration/startup/gpu/rpc_test.cpp
+++ b/libc/test/integration/startup/gpu/rpc_test.cpp
@@ -26,7 +26,6 @@ static void test_add_simple() {
         [&](LIBC_NAMESPACE::rpc::Buffer *buffer, uint32_t) {
           cnt = reinterpret_cast<uint64_t *>(buffer->data)[0];
         });
-    port.close();
   }
   ASSERT_TRUE(cnt == num_additions && "Incorrect sum");
 }
@@ -38,7 +37,6 @@ static void test_noop(uint8_t data) {
   port.send([=](LIBC_NAMESPACE::rpc::Buffer *buffer, uint32_t) {
     buffer->data[0] = data;
   });
-  port.close();
 }
 
 TEST_MAIN(int argc, char **argv, char **envp) {
diff --git a/llvm/tools/llvm-gpu-loader/server.h b/llvm/tools/llvm-gpu-loader/server.h
index da73cc007f5d..61606a12d7d3 100644
--- a/llvm/tools/llvm-gpu-loader/server.h
+++ b/llvm/tools/llvm-gpu-loader/server.h
@@ -47,8 +47,6 @@ inline uint32_t handle_server(rpc::Server &server, uint32_t index,
   if (status != rpc::RPC_SUCCESS)
     handle_error("Error handling RPC server");
 
-  port->close();
-
   return index;
 }
 
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index 234bd1061465..8bb60feab7b8 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -111,7 +111,6 @@ runServer(plugin::GenericDeviceTy &Device, void *Buffer,
   if (Status == rpc::RPC_UNHANDLED_OPCODE)
     Status = LIBC_NAMESPACE::shared::handle_libc_opcodes(*Port, NumLanes);
 
-  Port->close();
   return Status;
 }
 
diff --git a/openmp/device/src/Misc.cpp b/openmp/device/src/Misc.cpp
index 5d5a2a383f2b..158eac2f457c 100644
--- a/openmp/device/src/Misc.cpp
+++ b/openmp/device/src/Misc.cpp
@@ -131,7 +131,6 @@ unsigned long long __llvm_omp_host_call(void *fn, void *data, size_t size) {
   Port.recv([&](rpc::Buffer *Buffer, uint32_t) {
     Ret = static_cast<unsigned long long>(Buffer->data[0]);
   });
-  Port.close();
   return Ret;
 }
 }