[libc-commits] [libc] [libc][Docs] Update the GPU RPC documentation (PR #79069)
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Tue Jan 23 09:40:34 PST 2024
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/79069
>From 34332330ae18a3c6ec85401b07d56329e2733083 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Mon, 22 Jan 2024 16:38:07 -0600
Subject: [PATCH] [libc][Docs] Update the GPU RPC documentation
Summary:
This adds some more concrete information on the RPC interface. Hopefully
this is intelligable and provides some useful examples.
---
libc/docs/gpu/rpc-diagram.svg | 1 +
libc/docs/gpu/rpc.rst | 296 +++++++++++++++++++++++++++++++++-
libc/docs/gpu/testing.rst | 6 +-
3 files changed, 296 insertions(+), 7 deletions(-)
create mode 100644 libc/docs/gpu/rpc-diagram.svg
diff --git a/libc/docs/gpu/rpc-diagram.svg b/libc/docs/gpu/rpc-diagram.svg
new file mode 100644
index 00000000000000..eb9b807e52004e
--- /dev/null
+++ b/libc/docs/gpu/rpc-diagram.svg
@@ -0,0 +1 @@
+<svg version="1.1" viewBox="0.0 0.0 960.0 540.0" fill="none" stroke="none" stroke-linecap="square" stroke-miterlimit="10" xmlns:xlink="http://www.w3.org/1999/xlink" xmlns="http://www.w3.org/2000/svg"><clipPath id="g2b223c31cff_0_0.0"><path d="m0 0l960.0 0l0 540.0l-960.0 0l0 -540.0z" clip-rule="nonzero"/></clipPath><g clip-path="url(#g2b223c31cff_0_0.0)"><path fill="#000000" fill-opacity="0.0" d="m0 0l960.0 0l0 540.0l-960.0 0z" fill-rule="evenodd"/><path fill="#b6d7a8" d="m271.76743 215.15164l109.921265 0l0 79.93701l-109.921265 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m271.76743 215.15164l109.921265 0l0 79.93701l-109.921265 0z" fill-rule="evenodd"/><path fill="#000000" d="m294.84525 262.04013l0 -11.624985l1.03125 0l2.875 5.6874847l2.921875 -5.7031097l0.984375 0l0 11.64061l-1.234375 0l0 -8.76561l-2.515625 4.6874847l-0.5 0l-2.34375 -4.6406097l0 8.718735l-1.21875 0zm12.5 0.203125q-1.453125 0 -2.234375 -0.65625q-0.78125 -0.671875 -0.78125 -1.6875q0 -0.6875 0.328125 -1.234375q0.328125 -0.5625 0.921875 -0.9375q0.59375 -0.390625 1.375 -0.59375q0.78125 -0.203125 1.65625 -0.203125q0.3125 0 0.640625 0.015625q0.328125 0 0.6875 0.03125q0.375 0.03125 0.765625 0.078125l0.03125 1.03125q-0.328125 -0.0625 -0.671875 -0.078125q-0.34375 -0.03125 -0.6875 -0.046875q-0.328125 -0.03125 -0.65625 -0.03125q-0.671875 0 -1.234375 0.125q-0.546875 0.109375 -0.96875 0.328125q-0.40625 0.21875 -0.640625 0.578125q-0.21875 0.359375 -0.21875 0.859375q0 0.390625 0.140625 0.671875q0.140625 0.265625 0.375 0.421875q0.25 0.15625 0.578125 0.234375q0.328125 0.0625 0.71875 0.0625q0.734375 0 1.296875 -0.21875q0.578125 -0.234375 0.96875 -0.6875q0.40625 -0.453125 0.609375 -1.09375q0.203125 -0.65625 0.203125 -1.5q0 -1.265625 -0.296875 -1.9687347q-0.296875 -0.71875 -0.890625 -1.015625q-0.578125 -0.296875 -1.5 -0.296875q-0.609375 0 -1.203125 0.234375q-0.59375 0.234375 -1.078125 0.71875l-0.65625 -0.859375q0.59375 -0.59375 1.375 -0.890625q0.78125 -0.3125 1.65625 -0.3125q0.84375 0 1.53125 0.21875q0.6875 0.203125 1.1875 0.6875q0.5 0.46875 0.765625 1.265625q0.265625 0.79685974 0.265625 2.0156097l0 4.53125l-1.453125 0l0 -1.3125q-0.234375 0.46875 -0.59375 0.765625q-0.34375 0.296875 -0.734375 0.46875q-0.390625 0.15625 -0.796875 0.21875q-0.40625 0.0625 -0.78125 0.0625zm7.25 -0.203125l0 -1.078125l2.1875 0l0 -6.3593597l-2.0625 0l0 -1.09375l3.40625 0l0 7.4531097l2.0 0l0 1.078125l-5.53125 0zm2.78125 -10.312485q-0.390625 0 -0.671875 -0.28125q-0.28125 -0.28125 -0.28125 -0.671875q0 -0.40625 0.265625 -0.6875q0.28125 -0.28125 0.6875 -0.28125q0.390625 0 0.671875 0.296875q0.296875 0.28125 0.296875 0.671875q0 0.390625 -0.296875 0.671875q-0.28125 0.28125 -0.671875 0.28125zm6.15625 10.312485l0 -1.078125l2.53125 0l0 -10.249985l-2.421875 0l0 -1.078125l3.78125 0l0 11.32811l2.5 0l0 1.078125l-6.390625 0zm12.828125 0.203125q-0.53125 0 -1.0 -0.15625q-0.46875 -0.15625 -0.84375 -0.453125q-0.375 -0.296875 -0.6875 -0.71875l-0.453125 1.125l-0.859375 0l0 -12.406235l1.34375 0l0.1875 0l0 0.125q-0.125 0.125 -0.15625 0.25q-0.015625 0.125 -0.015625 0.453125l0 4.359375q0.203125 -0.34375 0.484375 -0.609375q0.28125 -0.28125 0.609375 -0.484375q0.34375 -0.21875 0.71875 -0.3125q0.375 -0.109375 0.71875 -0.109375q0.78125 0 1.421875 0.28125q0.65625 0.265625 1.109375 0.8125q0.46875 0.546875 0.71875 1.375q0.265625 0.81248474 0.265625 1.9218597q0 1.140625 -0.3125 2.0q-0.296875 0.84375 -0.8125 1.421875q-0.5 0.5625 -1.140625 0.84375q-0.625 0.28125 -1.296875 0.28125l0 0zm-0.171875 -1.1875q0.453125 0 0.859375 -0.171875q0.421875 -0.171875 0.75 -0.546875q0.34375 -0.390625 0.53125 -1.015625q0.203125 -0.640625 0.203125 -1.546875q0 -0.84375 -0.171875 -1.46875q-0.15625 -0.62498474 -0.46875 -1.0312347q-0.3125 -0.421875 -0.75 -0.609375q-0.421875 -0.203125 -0.921875 -0.203125q-0.6875 0 -1.21875 0.421875q-0.53125 0.421875 -0.828125 1.1874847q-0.296875 0.75 -0.296875 1.734375q0 0.734375 0.125 1.328125q0.140625 0.59375 0.40625 1.03125q0.265625 0.421875 0.703125 0.65625q0.4375 0.234375 1.078125 0.234375l0 0zm9.203125 1.140625q-1.171875 0 -2.0625 -0.5625q-0.875 -0.578125 -1.375 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3124847q0.5 -1.0 1.375 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 0.99998474 0.5 2.3124847q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.74998474 -0.890625 -1.1718597q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.890625 1.1718597q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.328125 0.75 0.890625 1.1875q0.578125 0.421875 1.3125 0.421875zm5.53125 0.96875l3.0 -4.328125l-2.9375 -4.2031097l1.46875 0l2.21875 3.1562347l2.125 -3.1562347l1.390625 0l-2.765625 4.1406097l3.125 4.390625l-1.515625 0l-2.359375 -3.328125l-2.21875 3.328125l-1.53125 0z" fill-rule="nonzero"/><path fill="#ffe599" d="m716.0079 215.15164l109.921265 0l0 79.93701l-109.921265 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m716.0079 215.15164l109.921265 0l0 79.93701l-109.921265 0z" fill-rule="evenodd"/><path fill="#000000" d="m739.0857 262.04013l0 -11.624985l1.03125 0l2.875 5.6874847l2.921875 -5.7031097l0.984375 0l0 11.64061l-1.234375 0l0 -8.76561l-2.515625 4.6874847l-0.5 0l-2.34375 -4.6406097l0 8.718735l-1.21875 0zm12.5 0.203125q-1.453125 0 -2.234375 -0.65625q-0.78125 -0.671875 -0.78125 -1.6875q0 -0.6875 0.328125 -1.234375q0.328125 -0.5625 0.921875 -0.9375q0.59375 -0.390625 1.375 -0.59375q0.78125 -0.203125 1.65625 -0.203125q0.3125 0 0.640625 0.015625q0.328125 0 0.6875 0.03125q0.375 0.03125 0.765625 0.078125l0.03125 1.03125q-0.328125 -0.0625 -0.671875 -0.078125q-0.34375 -0.03125 -0.6875 -0.046875q-0.328125 -0.03125 -0.65625 -0.03125q-0.671875 0 -1.234375 0.125q-0.546875 0.109375 -0.96875 0.328125q-0.40625 0.21875 -0.640625 0.578125q-0.21875 0.359375 -0.21875 0.859375q0 0.390625 0.140625 0.671875q0.140625 0.265625 0.375 0.421875q0.25 0.15625 0.578125 0.234375q0.328125 0.0625 0.71875 0.0625q0.734375 0 1.296875 -0.21875q0.578125 -0.234375 0.96875 -0.6875q0.40625 -0.453125 0.609375 -1.09375q0.203125 -0.65625 0.203125 -1.5q0 -1.265625 -0.296875 -1.9687347q-0.296875 -0.71875 -0.890625 -1.015625q-0.578125 -0.296875 -1.5 -0.296875q-0.609375 0 -1.203125 0.234375q-0.59375 0.234375 -1.078125 0.71875l-0.65625 -0.859375q0.59375 -0.59375 1.375 -0.890625q0.78125 -0.3125 1.65625 -0.3125q0.84375 0 1.53125 0.21875q0.6875 0.203125 1.1875 0.6875q0.5 0.46875 0.765625 1.265625q0.265625 0.79685974 0.265625 2.0156097l0 4.53125l-1.453125 0l0 -1.3125q-0.234375 0.46875 -0.59375 0.765625q-0.34375 0.296875 -0.734375 0.46875q-0.390625 0.15625 -0.796875 0.21875q-0.40625 0.0625 -0.78125 0.0625zm7.25 -0.203125l0 -1.078125l2.1875 0l0 -6.3593597l-2.0625 0l0 -1.09375l3.40625 0l0 7.4531097l2.0 0l0 1.078125l-5.53125 0zm2.78125 -10.312485q-0.390625 0 -0.671875 -0.28125q-0.28125 -0.28125 -0.28125 -0.671875q0 -0.40625 0.265625 -0.6875q0.28125 -0.28125 0.6875 -0.28125q0.390625 0 0.671875 0.296875q0.296875 0.28125 0.296875 0.671875q0 0.390625 -0.296875 0.671875q-0.28125 0.28125 -0.671875 0.28125zm6.15625 10.312485l0 -1.078125l2.53125 0l0 -10.249985l-2.421875 0l0 -1.078125l3.78125 0l0 11.32811l2.5 0l0 1.078125l-6.390625 0zm12.828125 0.203125q-0.53125 0 -1.0 -0.15625q-0.46875 -0.15625 -0.84375 -0.453125q-0.375 -0.296875 -0.6875 -0.71875l-0.453125 1.125l-0.859375 0l0 -12.406235l1.34375 0l0.1875 0l0 0.125q-0.125 0.125 -0.15625 0.25q-0.015625 0.125 -0.015625 0.453125l0 4.359375q0.203125 -0.34375 0.484375 -0.609375q0.28125 -0.28125 0.609375 -0.484375q0.34375 -0.21875 0.71875 -0.3125q0.375 -0.109375 0.71875 -0.109375q0.78125 0 1.421875 0.28125q0.65625 0.265625 1.109375 0.8125q0.46875 0.546875 0.71875 1.375q0.265625 0.81248474 0.265625 1.9218597q0 1.140625 -0.3125 2.0q-0.296875 0.84375 -0.8125 1.421875q-0.5 0.5625 -1.140625 0.84375q-0.625 0.28125 -1.296875 0.28125l0 0zm-0.171875 -1.1875q0.453125 0 0.859375 -0.171875q0.421875 -0.171875 0.75 -0.546875q0.34375 -0.390625 0.53125 -1.015625q0.203125 -0.640625 0.203125 -1.546875q0 -0.84375 -0.171875 -1.46875q-0.15625 -0.62498474 -0.46875 -1.0312347q-0.3125 -0.421875 -0.75 -0.609375q-0.421875 -0.203125 -0.921875 -0.203125q-0.6875 0 -1.21875 0.421875q-0.53125 0.421875 -0.828125 1.1874847q-0.296875 0.75 -0.296875 1.734375q0 0.734375 0.125 1.328125q0.140625 0.59375 0.40625 1.03125q0.265625 0.421875 0.703125 0.65625q0.4375 0.234375 1.078125 0.234375l0 0zm9.203125 1.140625q-1.171875 0 -2.0625 -0.5625q-0.875 -0.578125 -1.375 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3124847q0.5 -1.0 1.375 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 0.99998474 0.5 2.3124847q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.74998474 -0.890625 -1.1718597q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.890625 1.1718597q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.328125 0.75 0.890625 1.1875q0.578125 0.421875 1.3125 0.421875zm5.53125 0.96875l3.0 -4.328125l-2.9375 -4.2031097l1.46875 0l2.21875 3.1562347l2.125 -3.1562347l1.390625 0l-2.765625 4.1406097l3.125 4.390625l-1.515625 0l-2.359375 -3.328125l-2.21875 3.328125l-1.53125 0z" fill-rule="nonzero"/><path fill="#6aa84f" d="m417.9425 -3.874016l261.7953 0l0 91.71654l-261.7953 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.9425 -3.874016l261.7953 0l0 91.71654l-261.7953 0z" fill-rule="evenodd"/><path fill="#000000" d="m531.27765 48.90425l0 -11.625l3.734375 0q1.3125 0 2.109375 0.453125q0.8125 0.4375 1.1875 1.1875q0.375 0.734375 0.375 1.640625q0 0.671875 -0.203125 1.25q-0.203125 0.578125 -0.640625 1.03125q-0.4375 0.453125 -1.109375 0.703125q-0.671875 0.234375 -1.625 0.234375l-2.484375 0l0 5.125l-1.34375 0zm1.34375 -6.28125l2.34375 0q0.859375 0 1.375 -0.25q0.515625 -0.265625 0.75 -0.71875q0.25 -0.453125 0.25 -1.046875q0 -0.578125 -0.25 -1.046875q-0.234375 -0.46875 -0.75 -0.75q-0.515625 -0.296875 -1.359375 -0.296875l-2.359375 0l0 4.109375zm11.5625 6.4375q-1.171875 0 -2.0625 -0.5625q-0.875 -0.578125 -1.375 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3125q0.5 -1.0 1.375 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 1.0 0.5 2.3125q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.75 -0.890625 -1.171875q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.890625 1.171875q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.328125 0.75 0.890625 1.1875q0.578125 0.421875 1.3125 0.421875zm6.875 -7.5625l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875zm12.6875 8.65625q-0.828125 0 -1.34375 -0.265625q-0.515625 -0.28125 -0.796875 -0.875q-0.28125 -0.59375 -0.375 -1.484375q-0.078125 -0.890625 -0.046875 -2.125l0.234375 -6.078125l1.390625 -0.171875l0.1875 -0.015625l0.015625 0.125q-0.09375 0.125 -0.15625 0.265625q-0.0625 0.125 -0.09375 0.4375l-0.171875 2.109375l0.0625 0.34375l-0.140625 3.078125q-0.046875 1.421875 0.078125 2.203125q0.125 0.765625 0.453125 1.0625q0.34375 0.28125 0.90625 0.28125q0.625 0 1.09375 -0.234375q0.46875 -0.25 1.0 -0.625l0.40625 1.0625q-0.625 0.46875 -1.296875 0.6875q-0.671875 0.21875 -1.40625 0.21875l0 0zm-4.265625 -8.65625l6.125 0l0 1.109375l-6.125 0l0 -1.109375z" fill-rule="nonzero"/><path fill="#6aa84f" d="m417.95084 420.30414l261.79526 0l0 91.71655l-261.79526 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.95084 420.30414l261.79526 0l0 91.71655l-261.79526 0z" fill-rule="evenodd"/><path fill="#000000" d="m531.28595 473.0824l0 -11.625l3.734375 0q1.3125 0 2.109375 0.453125q0.8125 0.4375 1.1875 1.1875q0.375 0.734375 0.375 1.640625q0 0.671875 -0.203125 1.25q-0.203125 0.578125 -0.640625 1.03125q-0.4375 0.453125 -1.109375 0.703125q-0.671875 0.234375 -1.625 0.234375l-2.484375 0l0 5.125l-1.34375 0zm1.34375 -6.28125l2.34375 0q0.859375 0 1.375 -0.25q0.515625 -0.265625 0.75 -0.71875q0.25 -0.453125 0.25 -1.046875q0 -0.578125 -0.25 -1.046875q-0.234375 -0.46875 -0.75 -0.75q-0.515625 -0.296875 -1.359375 -0.296875l-2.359375 0l0 4.109375zm11.5625 6.4375q-1.171875 0 -2.0625 -0.5625q-0.875 -0.578125 -1.375 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3125q0.5 -1.0 1.375 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 1.0 0.5 2.3125q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.75 -0.890625 -1.171875q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.890625 1.171875q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.328125 0.75 0.890625 1.1875q0.578125 0.421875 1.3125 0.421875zm6.875 -7.5625l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875zm12.6875 8.65625q-0.828125 0 -1.34375 -0.265625q-0.515625 -0.28125 -0.796875 -0.875q-0.28125 -0.59375 -0.375 -1.484375q-0.078125 -0.890625 -0.046875 -2.125l0.234375 -6.078125l1.390625 -0.171875l0.1875 -0.015625l0.015625 0.125q-0.09375 0.125 -0.15625 0.265625q-0.0625 0.125 -0.09375 0.4375l-0.171875 2.109375l0.0625 0.34375l-0.140625 3.078125q-0.046875 1.421875 0.078125 2.203125q0.125 0.765625 0.453125 1.0625q0.34375 0.28125 0.90625 0.28125q0.625 0 1.09375 -0.234375q0.46875 -0.25 1.0 -0.625l0.40625 1.0625q-0.625 0.46875 -1.296875 0.6875q-0.671875 0.21875 -1.40625 0.21875l0 0zm-4.265625 -8.65625l6.125 0l0 1.109375l-6.125 0l0 -1.109375z" fill-rule="nonzero"/><path fill="#000000" fill-opacity="0.0" d="m417.95084 466.1624c-45.606293 0 -91.212585 -85.543335 -91.212585 -171.08664" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.95084 466.1624c-22.803162 0 -45.606293 -21.385834 -62.70868 -53.46457c-8.551178 -16.039368 -15.677155 -34.751984 -20.665344 -54.801178c-2.4940796 -10.024628 -4.4537354 -20.383392 -5.789856 -30.90921c-0.6680603 -5.2629395 -1.1802368 -10.567627 -1.5254211 -15.893188c-0.08627319 -1.3313904 -0.1621399 -2.664093 -0.22738647 -3.9977722l-0.0011291504 -0.024353027" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m330.33548 306.99103l-3.5254211 -8.99231l-3.0794983 9.154602z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m326.72806 215.15164c0 -86.58269 45.606293 -173.16536 91.212585 -173.16536" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m326.72806 215.15164c0 -43.29135 11.401581 -86.58269 28.503937 -119.051186c8.551178 -16.234253 18.527557 -29.762794 29.216522 -39.232777c5.344513 -4.734989 10.867126 -8.455341 16.478851 -10.99194c1.4029236 -0.63415146 2.811432 -1.1943207 4.2240906 -1.6778603c0.1765747 -0.06044388 0.35327148 -0.11968994 0.52993774 -0.17773438l0.4109192 -0.13211441" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m406.61584 47.149746l8.437958 -4.7001038l-9.485016 -1.8233299z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m770.9685 295.08865c0 85.543304 -45.606323 171.08661 -91.212585 171.08661" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m770.9685 295.08865c0 42.771637 -11.40155 85.543274 -28.503967 117.62201c-8.551147 16.039368 -18.527527 29.405518 -29.216492 38.76181c-5.3444824 4.6781616 -10.867126 8.353851 -16.478882 10.859985c-1.402832 0.6265869 -2.8114014 1.1800232 -4.22406 1.6577454c-0.1765747 0.059692383 -0.35327148 0.1182251 -0.5299072 0.17556763l-0.40753174 0.12945557" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m691.09015 461.03253l-8.446533 4.6846313l9.481628 1.8407288z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m679.7378 41.984253c45.62207 0 91.24408 86.58267 91.24408 173.16536" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m679.7378 41.984253c22.811035 0 45.62201 21.645668 62.730286 54.114174c8.554138 16.234253 15.682617 35.174217 20.672485 55.467033c2.494934 10.146393 4.4553223 20.631027 5.79187 31.284744c0.6682739 5.326874 1.1806641 10.696014 1.526001 16.086288c0.086242676 1.3475647 0.16210938 2.696457 0.22741699 4.046341l0.007873535 0.17024231" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m767.39124 203.23239l3.5204468 8.994247l3.0845947 -9.152893z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m534.39343 151.2861l0 -63.46457" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m534.39343 151.2861l0 -51.46457" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m537.6969 99.82153l-3.3034668 -9.076195l-3.3034668 9.076195z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m563.27734 87.80994l0 63.46457" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m563.27734 87.80994l0 51.46457" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m559.9739 139.2745l3.3034668 9.076187l3.3034668 -9.076187z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m563.28564 420.307l0 -63.46457" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m563.2857 420.307l0 -51.46457" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m566.5892 368.84244l-3.3034668 -9.076202l-3.3034668 9.076202z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m534.40173 356.83084l0 63.46457" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m534.40173 356.83084l0 51.46457" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m531.09827 408.2954l3.3034668 9.076172l3.3034668 -9.076172z" fill-rule="evenodd"/><path fill="#ea9999" d="m417.9342 151.2861l131.9685 0l0 54.929123l-131.9685 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.9342 151.2861l131.9685 0l0 54.929123l-131.9685 0z" fill-rule="evenodd"/><path fill="#000000" d="m460.60596 185.8738q-1.328125 0 -2.265625 -0.6875q-0.9375 -0.6875 -1.421875 -2.015625q-0.484375 -1.34375 -0.484375 -3.25q0 -1.84375 0.484375 -3.171875q0.484375 -1.34375 1.421875 -2.0625q0.9375 -0.71875 2.265625 -0.71875q1.328125 0 2.25 0.71875q0.9375 0.71875 1.421875 2.0625q0.484375 1.328125 0.484375 3.171875q0 1.90625 -0.484375 3.25q-0.484375 1.328125 -1.421875 2.015625q-0.921875 0.6875 -2.25 0.6875zm0 -1.140625q0.890625 0 1.515625 -0.546875q0.625 -0.5625 0.953125 -1.640625q0.34375 -1.078125 0.34375 -2.625q0 -1.46875 -0.34375 -2.53125q-0.328125 -1.078125 -0.953125 -1.671875q-0.625 -0.609375 -1.515625 -0.609375q-0.890625 0 -1.53125 0.609375q-0.625 0.59375 -0.96875 1.671875q-0.328125 1.0625 -0.328125 2.53125q0 1.546875 0.328125 2.625q0.34375 1.078125 0.96875 1.640625q0.640625 0.546875 1.53125 0.546875zm5.78125 4.046875l0 -11.640625l1.359375 0l0 1.6875q0.203125 -0.65625 0.578125 -1.0625q0.390625 -0.421875 0.90625 -0.609375q0.515625 -0.203125 1.109375 -0.203125q0.703125 0 1.34375 0.28125q0.65625 0.265625 1.140625 0.8125q0.5 0.53125 0.78125 1.359375q0.296875 0.8125 0.296875 1.9375q0 1.125 -0.3125 1.984375q-0.3125 0.84375 -0.828125 1.421875q-0.5 0.5625 -1.15625 0.859375q-0.640625 0.28125 -1.3125 0.28125q-0.5 0 -0.984375 -0.15625q-0.46875 -0.15625 -0.859375 -0.453125q-0.390625 -0.3125 -0.6875 -0.734375l0 4.234375l-1.375 0zm3.6875 -4.078125q0.453125 0 0.875 -0.15625q0.4375 -0.15625 0.78125 -0.53125q0.359375 -0.375 0.5625 -1.015625q0.203125 -0.65625 0.203125 -1.59375q0 -1.03125 -0.265625 -1.75q-0.25 -0.71875 -0.75 -1.109375q-0.484375 -0.390625 -1.203125 -0.4375q-0.515625 -0.03125 -0.984375 0.15625q-0.453125 0.1875 -0.8125 0.625q-0.34375 0.421875 -0.546875 1.125q-0.1875 0.6875 -0.1875 1.6875q0.015625 0.59375 0.125 1.078125q0.109375 0.484375 0.3125 0.859375q0.21875 0.359375 0.484375 0.59375q0.28125 0.21875 0.640625 0.34375q0.359375 0.125 0.765625 0.125zm9.828125 1.171875q-0.921875 0 -1.71875 -0.328125q-0.78125 -0.34375 -1.359375 -0.9375q-0.5625 -0.59375 -0.875 -1.40625q-0.3125 -0.8125 -0.3125 -1.78125q0 -0.953125 0.3125 -1.765625q0.3125 -0.8125 0.875 -1.40625q0.578125 -0.59375 1.359375 -0.921875q0.796875 -0.34375 1.703125 -0.34375q1.046875 0 1.875 0.421875q0.828125 0.40625 1.359375 1.15625l-0.84375 0.828125l-0.125 0.125l-0.09375 -0.09375q0 -0.15625 -0.078125 -0.28125q-0.0625 -0.125 -0.28125 -0.359375q-0.40625 -0.359375 -0.875 -0.515625q-0.46875 -0.15625 -1.078125 -0.15625q-0.546875 0 -1.03125 0.234375q-0.484375 0.21875 -0.875 0.640625q-0.375 0.421875 -0.609375 1.015625q-0.21875 0.59375 -0.21875 1.328125q0 0.71875 0.21875 1.328125q0.234375 0.609375 0.625 1.0625q0.390625 0.4375 0.9375 0.703125q0.546875 0.25 1.1875 0.25q0.4375 0 0.828125 -0.125q0.390625 -0.125 0.734375 -0.359375q0.359375 -0.234375 0.65625 -0.5625l0.78125 0.90625q-0.625 0.671875 -1.40625 1.015625q-0.765625 0.328125 -1.671875 0.328125l0 0zm8.6875 -0.046875q-1.171875 0 -2.0625 -0.5625q-0.875 -0.578125 -1.375 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3125q0.5 -1.0 1.375 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 1.0 0.5 2.3125q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.75 -0.890625 -1.171875q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.890625 1.171875q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.328125 0.75 0.890625 1.1875q0.578125 0.421875 1.3125 0.421875zm9.015625 1.15625q-0.640625 0 -1.265625 -0.28125q-0.625 -0.28125 -1.140625 -0.84375q-0.515625 -0.5625 -0.8125 -1.40625q-0.296875 -0.859375 -0.296875 -2.0q0 -1.171875 0.3125 -2.0q0.3125 -0.84375 0.828125 -1.375q0.53125 -0.53125 1.15625 -0.78125q0.640625 -0.25 1.296875 -0.25q0.59375 0 1.0625 0.1875q0.484375 0.1875 0.84375 0.515625q0.375 0.328125 0.578125 0.765625l0 -5.125l1.234375 0l0.1875 0l0 0.125q-0.109375 0.125 -0.140625 0.25q-0.03125 0.125 -0.03125 0.453125l0.015625 10.25q0 0.34375 0.03125 0.671875q0.03125 0.328125 0.15625 0.65625l-1.328125 0q-0.09375 -0.21875 -0.125 -0.4375q-0.03125 -0.234375 -0.046875 -0.4375q-0.015625 -0.21875 -0.015625 -0.453125q-0.25 0.453125 -0.625 0.796875q-0.375 0.34375 -0.84375 0.53125q-0.46875 0.1875 -1.03125 0.1875l0 0zm0.21875 -1.171875q0.609375 0 1.03125 -0.25q0.4375 -0.265625 0.6875 -0.71875q0.265625 -0.453125 0.390625 -1.0625q0.125 -0.609375 0.125 -1.296875q0 -0.953125 -0.25 -1.703125q-0.234375 -0.75 -0.765625 -1.1875q-0.515625 -0.453125 -1.34375 -0.453125q-0.546875 0 -1.0 0.21875q-0.4375 0.21875 -0.734375 0.625q-0.28125 0.40625 -0.4375 0.984375q-0.140625 0.5625 -0.140625 1.28125q0 1.0625 0.28125 1.859375q0.28125 0.796875 0.8125 1.25q0.546875 0.453125 1.34375 0.453125l0 0zm9.828125 1.171875q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625z" fill-rule="nonzero"/><path fill="#a4c2f4" d="m134.06955 -3.874016l101.44882 0l0 150.61418l-101.44882 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m134.06955 -3.874016l101.44882 0l0 150.61418l-101.44882 0z" fill-rule="evenodd"/><path fill="#000000" d="m161.41896 78.540565q-1.140625 0 -2.046875 -0.375q-0.90625 -0.375 -1.59375 -1.09375l0.671875 -1.1875l0.078125 -0.15625l0.125 0.0625q0.046875 0.140625 0.09375 0.296875q0.0625 0.15625 0.265625 0.375q0.484375 0.421875 1.125 0.671875q0.640625 0.234375 1.46875 0.234375q0.515625 0 0.953125 -0.15625q0.4375 -0.171875 0.765625 -0.4375q0.328125 -0.265625 0.5 -0.625q0.1875 -0.375 0.1875 -0.78125q0 -0.3125 -0.0625 -0.578125q-0.0625 -0.265625 -0.25 -0.46875q-0.171875 -0.21875 -0.46875 -0.4375q-0.296875 -0.21875 -0.765625 -0.453125q-0.46875 -0.25 -1.125 -0.53125q-1.203125 -0.453125 -1.90625 -1.0q-0.703125 -0.5625 -1.0 -1.171875q-0.28125 -0.625 -0.28125 -1.28125q0 -0.84375 0.421875 -1.484375q0.4375 -0.65625 1.21875 -1.015625q0.796875 -0.359375 1.90625 -0.359375q0.625 0 1.21875 0.171875q0.59375 0.171875 1.09375 0.5q0.515625 0.3125 0.90625 0.78125l-0.71875 0.96875l-0.109375 0.140625l-0.125 -0.0625q-0.015625 -0.15625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.453125 -0.4375 -0.953125 -0.59375q-0.5 -0.15625 -1.1875 -0.15625q-0.46875 0 -0.84375 0.140625q-0.359375 0.125 -0.625 0.375q-0.25 0.234375 -0.390625 0.546875q-0.140625 0.296875 -0.140625 0.625q0 0.328125 0.0625 0.59375q0.078125 0.25 0.265625 0.5q0.1875 0.234375 0.515625 0.484375q0.328125 0.234375 0.8125 0.484375q0.5 0.234375 1.21875 0.515625q0.859375 0.359375 1.421875 0.703125q0.578125 0.34375 0.890625 0.734375q0.328125 0.375 0.46875 0.828125q0.140625 0.4375 0.140625 0.984375q0 0.796875 -0.40625 1.578125q-0.390625 0.765625 -1.25 1.28125q-0.859375 0.5 -2.21875 0.5l0 0zm9.796875 0q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625zm9.296875 -3.421875l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875zm11.234375 8.53125l-3.375 -8.53125l1.34375 0l2.75 7.0l-0.25 0l1.5 -3.46875q0.453125 -1.03125 0.71875 -1.90625q0.28125 -0.875 0.40625 -1.625l1.265625 0q-0.140625 0.78125 -0.46875 1.71875q-0.3125 0.9375 -0.78125 2.015625l-2.0625 4.796875l-1.046875 0zm10.28125 0.1875q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625zm9.296875 -3.421875l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875z" fill-rule="nonzero"/><path fill="#a4c2f4" d="m134.06955 356.83084l101.44882 0l0 155.14963l-101.44882 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m134.06955 356.83084l101.44882 0l0 155.14963l-101.44882 0z" fill-rule="evenodd"/><path fill="#000000" d="m162.12209 441.51315q-0.953125 0 -1.78125 -0.34375q-0.8125 -0.34375 -1.4375 -1.0625q-0.625 -0.734375 -0.984375 -1.859375q-0.34375 -1.125 -0.34375 -2.671875q0 -1.359375 0.25 -2.375q0.265625 -1.015625 0.71875 -1.703125q0.453125 -0.6875 1.03125 -1.09375q0.578125 -0.421875 1.203125 -0.609375q0.640625 -0.1875 1.28125 -0.1875q1.265625 0 2.15625 0.625q0.890625 0.625 1.34375 1.671875l-1.125 0.515625l-0.15625 0.078125l-0.0625 -0.125q0.0625 -0.15625 0 -0.265625q-0.0625 -0.125 -0.203125 -0.390625q-0.46875 -0.640625 -0.96875 -0.84375q-0.484375 -0.21875 -1.03125 -0.21875q-0.71875 0 -1.3125 0.328125q-0.59375 0.328125 -1.015625 0.953125q-0.40625 0.625 -0.625 1.515625q-0.21875 0.890625 -0.21875 2.03125q0 1.109375 0.25 2.0q0.25 0.890625 0.6875 1.546875q0.453125 0.65625 1.03125 1.015625q0.59375 0.34375 1.265625 0.34375q0.46875 0 0.90625 -0.15625q0.4375 -0.171875 0.796875 -0.5q0.375 -0.328125 0.640625 -0.75l0.96875 0.625q-0.5625 0.9375 -1.34375 1.421875q-0.78125 0.484375 -1.921875 0.484375l0 0zm5.484375 -0.1875l0 -1.078125l2.53125 0l0 -10.25l-2.421875 0l0 -1.078125l3.78125 0l0 11.328125l2.5 0l0 1.078125l-6.390625 0zm9.71875 0l0 -1.078125l2.1875 0l0 -6.359375l-2.0625 0l0 -1.09375l3.40625 0l0 7.453125l2.0 0l0 1.078125l-5.53125 0zm2.78125 -10.3125q-0.390625 0 -0.671875 -0.28125q-0.28125 -0.28125 -0.28125 -0.671875q0 -0.40625 0.265625 -0.6875q0.28125 -0.28125 0.6875 -0.28125q0.390625 0 0.671875 0.296875q0.296875 0.28125 0.296875 0.671875q0 0.390625 -0.296875 0.671875q-0.28125 0.28125 -0.671875 0.28125zm9.765625 10.5q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625zm8.421875 5.109375l0 -8.53125l1.3125 0l0 1.515625q0.328125 -0.515625 0.765625 -0.890625q0.453125 -0.390625 1.0 -0.609375q0.546875 -0.21875 1.125 -0.21875q0.703125 0 1.265625 0.359375q0.5625 0.359375 0.890625 1.140625q0.34375 0.78125 0.34375 2.0625l0 5.171875l-1.3125 0l0 -5.125q0 -0.90625 -0.203125 -1.4375q-0.203125 -0.53125 -0.5625 -0.765625q-0.359375 -0.234375 -0.796875 -0.234375q-0.4375 0 -0.875 0.1875q-0.4375 0.1875 -0.8125 0.53125q-0.375 0.34375 -0.609375 0.84375q-0.21875 0.484375 -0.21875 1.09375l0 4.90625l-1.3125 0zm13.5625 0.125q-0.828125 0 -1.34375 -0.265625q-0.515625 -0.28125 -0.796875 -0.875q-0.28125 -0.59375 -0.375 -1.484375q-0.078125 -0.890625 -0.046875 -2.125l0.234375 -6.078125l1.390625 -0.171875l0.1875 -0.015625l0.015625 0.125q-0.09375 0.125 -0.15625 0.265625q-0.0625 0.125 -0.09375 0.4375l-0.171875 2.109375l0.0625 0.34375l-0.140625 3.078125q-0.046875 1.421875 0.078125 2.203125q0.125 0.765625 0.453125 1.0625q0.34375 0.28125 0.90625 0.28125q0.625 0 1.09375 -0.234375q0.46875 -0.25 1.0 -0.625l0.40625 1.0625q-0.625 0.46875 -1.296875 0.6875q-0.671875 0.21875 -1.40625 0.21875l0 0zm-4.265625 -8.65625l6.125 0l0 1.109375l-6.125 0l0 -1.109375z" fill-rule="nonzero"/><path fill="#e06666" d="m271.76746 -3.874016l77.700775 0l0 91.71654l-77.700775 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m271.76746 -3.874016l77.700775 0l0 91.71654l-77.700775 0z" fill-rule="evenodd"/><path fill="#000000" d="m293.19595 37.27925l1.34375 0l0.1875 0l0 0.125q-0.109375 0.125 -0.140625 0.25q-0.03125 0.125 -0.03125 0.453125l0 9.671875l5.59375 0l0 1.125l-6.953125 0l0 -11.625zm12.7656555 11.78125q-1.171875 0 -2.0625 -0.5625q-0.8750305 -0.578125 -1.3750305 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3125q0.5 -1.0 1.3750305 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 1.0 0.5 2.3125q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.75 -0.890625 -1.171875q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.8906555 1.171875q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.32815552 0.75 0.8906555 1.1875q0.578125 0.421875 1.3125 0.421875zm9.96875 1.171875q-0.921875 0 -1.71875 -0.328125q-0.78125 -0.34375 -1.359375 -0.9375q-0.5625 -0.59375 -0.875 -1.40625q-0.3125 -0.8125 -0.3125 -1.78125q0 -0.953125 0.3125 -1.765625q0.3125 -0.8125 0.875 -1.40625q0.578125 -0.59375 1.359375 -0.921875q0.796875 -0.34375 1.703125 -0.34375q1.046875 0 1.875 0.421875q0.828125 0.40625 1.359375 1.15625l-0.84375 0.828125l-0.125 0.125l-0.09375 -0.09375q0 -0.15625 -0.078125 -0.28125q-0.0625 -0.125 -0.28125 -0.359375q-0.40625 -0.359375 -0.875 -0.515625q-0.46875 -0.15625 -1.078125 -0.15625q-0.546875 0 -1.03125 0.234375q-0.484375 0.21875 -0.875 0.640625q-0.375 0.421875 -0.609375 1.015625q-0.21875 0.59375 -0.21875 1.328125q0 0.71875 0.21875 1.328125q0.234375 0.609375 0.625 1.0625q0.390625 0.4375 0.9375 0.703125q0.546875 0.25 1.1875 0.25q0.4375 0 0.828125 -0.125q0.390625 -0.125 0.734375 -0.359375q0.359375 -0.234375 0.65625 -0.5625l0.78125 0.90625q-0.625 0.671875 -1.40625 1.015625q-0.765625 0.328125 -1.671875 0.328125l0 0zm7.734375 -4.46875l-1.125 1.03125l0 3.234375l-1.359375 0l0 -12.40625l1.34375 0l0.1875 0l0 0.125q-0.109375 0.125 -0.140625 0.25q-0.03125 0.125 -0.03125 0.453125l0 6.96875l4.265625 -3.953125q0.125 0.015625 0.234375 0.046875q0.109375 0.015625 0.21875 0.03125q0.125 0 0.265625 0.015625q0.140625 0 0.28125 0l0.484375 0l-3.59375 3.375l4.1875 5.09375l-1.734375 0.078125l-3.484375 -4.34375z" fill-rule="nonzero"/><path fill="#e06666" d="m271.76746 421.31134l77.700775 0l0 91.71655l-77.700775 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m271.76746 421.31134l77.700775 0l0 91.71655l-77.700775 0z" fill-rule="evenodd"/><path fill="#000000" d="m293.19595 462.46463l1.34375 0l0.1875 0l0 0.125q-0.109375 0.125 -0.140625 0.25q-0.03125 0.125 -0.03125 0.453125l0 9.671875l5.59375 0l0 1.125l-6.953125 0l0 -11.625zm12.7656555 11.78125q-1.171875 0 -2.0625 -0.5625q-0.8750305 -0.578125 -1.3750305 -1.578125q-0.5 -1.0 -0.5 -2.28125q0 -1.3125 0.5 -2.3125q0.5 -1.0 1.3750305 -1.5625q0.890625 -0.578125 2.0625 -0.578125q1.15625 0 2.03125 0.578125q0.890625 0.5625 1.390625 1.5625q0.5 1.0 0.5 2.3125q0 1.28125 -0.5 2.28125q-0.5 1.0 -1.390625 1.578125q-0.875 0.5625 -2.03125 0.5625zm0 -1.125q0.734375 0 1.296875 -0.421875q0.5625 -0.4375 0.890625 -1.1875q0.328125 -0.75 0.328125 -1.71875q0 -0.96875 -0.328125 -1.703125q-0.328125 -0.75 -0.890625 -1.171875q-0.5625 -0.421875 -1.296875 -0.421875q-0.734375 0 -1.3125 0.421875q-0.5625 0.421875 -0.8906555 1.171875q-0.328125 0.734375 -0.328125 1.703125q0 0.96875 0.328125 1.71875q0.32815552 0.75 0.8906555 1.1875q0.578125 0.421875 1.3125 0.421875zm9.96875 1.171875q-0.921875 0 -1.71875 -0.328125q-0.78125 -0.34375 -1.359375 -0.9375q-0.5625 -0.59375 -0.875 -1.40625q-0.3125 -0.8125 -0.3125 -1.78125q0 -0.953125 0.3125 -1.765625q0.3125 -0.8125 0.875 -1.40625q0.578125 -0.59375 1.359375 -0.921875q0.796875 -0.34375 1.703125 -0.34375q1.046875 0 1.875 0.421875q0.828125 0.40625 1.359375 1.15625l-0.84375 0.828125l-0.125 0.125l-0.09375 -0.09375q0 -0.15625 -0.078125 -0.28125q-0.0625 -0.125 -0.28125 -0.359375q-0.40625 -0.359375 -0.875 -0.515625q-0.46875 -0.15625 -1.078125 -0.15625q-0.546875 0 -1.03125 0.234375q-0.484375 0.21875 -0.875 0.640625q-0.375 0.421875 -0.609375 1.015625q-0.21875 0.59375 -0.21875 1.328125q0 0.71875 0.21875 1.328125q0.234375 0.609375 0.625 1.0625q0.390625 0.4375 0.9375 0.703125q0.546875 0.25 1.1875 0.25q0.4375 0 0.828125 -0.125q0.390625 -0.125 0.734375 -0.359375q0.359375 -0.234375 0.65625 -0.5625l0.78125 0.90625q-0.625 0.671875 -1.40625 1.015625q-0.765625 0.328125 -1.671875 0.328125l0 0zm7.734375 -4.46875l-1.125 1.03125l0 3.234375l-1.359375 0l0 -12.40625l1.34375 0l0.1875 0l0 0.125q-0.109375 0.125 -0.140625 0.25q-0.03125 0.125 -0.03125 0.453125l0 6.96875l4.265625 -3.953125q0.125 0.015625 0.234375 0.046875q0.109375 0.015625 0.21875 0.03125q0.125 0 0.265625 0.015625q0.140625 0 0.28125 0l0.484375 0l-3.59375 3.375l4.1875 5.09375l-1.734375 0.078125l-3.484375 -4.34375z" fill-rule="nonzero"/><path fill="#000000" fill-opacity="0.0" d="m349.46823 41.984253l68.47244 0" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m349.46823 41.984253l56.472412 0" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m405.94064 45.287716l9.076202 -3.303463l-9.076202 -3.3034668z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m349.46823 467.16962l68.47244 -1.0078735" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m349.46823 467.16962l56.473724 -0.8312378" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m405.99057 469.64148l9.026611 -3.4367065l-9.12384 -3.1695251z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m237.31076 41.984253l34.456696 0" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m237.31076 41.984253l22.456696 0" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m259.76746 45.287716l9.076202 -3.303463l-9.076202 -3.3034668z" fill-rule="evenodd"/><path fill="#000000" fill-opacity="0.0" d="m235.10603 467.16962l36.661423 0" fill-rule="evenodd"/><path stroke="#ffab40" stroke-width="2.0" stroke-linejoin="round" stroke-linecap="butt" d="m235.10603 467.16962l24.661423 0" fill-rule="evenodd"/><path fill="#ffab40" stroke="#ffab40" stroke-width="2.0" stroke-linecap="butt" d="m259.76746 470.47308l9.076202 -3.3034668l-9.076202 -3.3034668z" fill-rule="evenodd"/><path fill="#0097a7" d="m417.95084 206.67488l261.79526 0l0 37.669296l-261.79526 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.95084 206.67488l261.79526 0l0 37.669296l-261.79526 0z" fill-rule="evenodd"/><path fill="#000000" d="m521.7547 232.42952l0 -11.625l3.3125 0q1.0 0 1.71875 0.203125q0.734375 0.203125 1.203125 0.578125q0.484375 0.375 0.703125 0.90625q0.234375 0.515625 0.234375 1.125q0 0.546875 -0.203125 1.078125q-0.1875 0.515625 -0.578125 0.90625q-0.375 0.390625 -0.9375 0.59375q0.484375 0.171875 0.875 0.46875q0.40625 0.296875 0.6875 0.703125q0.296875 0.390625 0.453125 0.859375q0.171875 0.453125 0.171875 0.9375q0 0.734375 -0.265625 1.34375q-0.265625 0.59375 -0.796875 1.03125q-0.53125 0.421875 -1.359375 0.65625q-0.8125 0.234375 -1.9375 0.234375l-3.28125 0zm1.296875 -1.140625l2.21875 0q0.8125 0 1.421875 -0.203125q0.625 -0.21875 0.96875 -0.703125q0.34375 -0.484375 0.34375 -1.296875q0 -0.6875 -0.296875 -1.21875q-0.296875 -0.53125 -0.96875 -0.828125q-0.65625 -0.3125 -1.796875 -0.3125l-1.890625 0l0 4.5625zm0 -5.671875l1.828125 0q1.015625 0 1.609375 -0.234375q0.609375 -0.25 0.859375 -0.671875q0.265625 -0.421875 0.265625 -0.9375q0 -0.390625 -0.140625 -0.71875q-0.140625 -0.34375 -0.4375 -0.59375q-0.296875 -0.25 -0.78125 -0.390625q-0.484375 -0.15625 -1.171875 -0.15625l-2.03125 0l0 3.703125zm11.265625 7.015625q-0.5 0 -0.953125 -0.15625q-0.453125 -0.140625 -0.8125 -0.4375q-0.359375 -0.3125 -0.625 -0.78125q-0.265625 -0.46875 -0.40625 -1.125q-0.140625 -0.65625 -0.140625 -1.46875l0.03125 -4.765625l1.296875 0l0 4.765625q0 0.796875 0.140625 1.359375q0.15625 0.5625 0.421875 0.90625q0.28125 0.328125 0.625 0.484375q0.359375 0.140625 0.75 0.140625q0.453125 0 0.859375 -0.171875q0.421875 -0.1875 0.734375 -0.53125q0.328125 -0.359375 0.515625 -0.921875q0.1875 -0.5625 0.1875 -1.328125l0 -4.703125l1.3125 0l0 7.203125q0 0.34375 0.015625 0.671875q0.03125 0.328125 0.171875 0.65625l-1.296875 0q-0.09375 -0.21875 -0.125 -0.4375q-0.03125 -0.21875 -0.046875 -0.4375q-0.015625 -0.21875 -0.015625 -0.4375q-0.203125 0.359375 -0.484375 0.640625q-0.28125 0.28125 -0.625 0.484375q-0.34375 0.1875 -0.734375 0.28125q-0.375 0.109375 -0.796875 0.109375zm8.21875 -7.671875l-1.84375 0l0 -1.0625l1.84375 0l0 -0.703125q0 -1.140625 0.40625 -1.859375q0.40625 -0.734375 1.109375 -1.078125q0.71875 -0.34375 1.625 -0.34375q0.796875 0 1.5 0.296875q0.703125 0.28125 1.171875 0.875l-0.71875 0.984375l-0.109375 0.15625l-0.109375 -0.078125q-0.015625 -0.140625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.3125 -0.265625 -0.609375 -0.375q-0.296875 -0.109375 -0.765625 -0.109375q-0.53125 0 -0.96875 0.21875q-0.421875 0.21875 -0.671875 0.6875q-0.234375 0.46875 -0.234375 1.296875l0 0.703125l2.796875 0l0 1.0625l-2.796875 0l0 7.46875l-1.296875 0l0 -7.46875zm9.328125 0l-1.84375 0l0 -1.0625l1.84375 0l0 -0.703125q0 -1.140625 0.40625 -1.859375q0.40625 -0.734375 1.109375 -1.078125q0.71875 -0.34375 1.625 -0.34375q0.796875 0 1.5 0.296875q0.703125 0.28125 1.171875 0.875l-0.71875 0.984375l-0.109375 0.15625l-0.109375 -0.078125q-0.015625 -0.140625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.3125 -0.265625 -0.609375 -0.375q-0.296875 -0.109375 -0.765625 -0.109375q-0.53125 0 -0.96875 0.21875q-0.421875 0.21875 -0.671875 0.6875q-0.234375 0.46875 -0.234375 1.296875l0 0.703125l2.796875 0l0 1.0625l-2.796875 0l0 7.46875l-1.296875 0l0 -7.46875zm11.390625 7.65625q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625zm9.296875 -3.421875l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875z" fill-rule="nonzero"/><path fill="#0097a7" d="m417.95084 244.3293l261.79526 0l0 37.669296l-261.79526 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.95084 244.3293l261.79526 0l0 37.669296l-261.79526 0z" fill-rule="evenodd"/><path fill="#000000" d="m521.7547 270.08395l0 -11.625l3.3125 0q1.0 0 1.71875 0.203125q0.734375 0.203125 1.203125 0.578125q0.484375 0.375 0.703125 0.90625q0.234375 0.515625 0.234375 1.125q0 0.546875 -0.203125 1.078125q-0.1875 0.515625 -0.578125 0.90625q-0.375 0.390625 -0.9375 0.59375q0.484375 0.171875 0.875 0.46875q0.40625 0.296875 0.6875 0.703125q0.296875 0.390625 0.453125 0.859375q0.171875 0.453125 0.171875 0.9375q0 0.734375 -0.265625 1.34375q-0.265625 0.59375 -0.796875 1.03125q-0.53125 0.421875 -1.359375 0.65625q-0.8125 0.234375 -1.9375 0.234375l-3.28125 0zm1.296875 -1.140625l2.21875 0q0.8125 0 1.421875 -0.203125q0.625 -0.21875 0.96875 -0.703125q0.34375 -0.484375 0.34375 -1.296875q0 -0.6875 -0.296875 -1.21875q-0.296875 -0.53125 -0.96875 -0.828125q-0.65625 -0.3125 -1.796875 -0.3125l-1.890625 0l0 4.5625zm0 -5.671875l1.828125 0q1.015625 0 1.609375 -0.234375q0.609375 -0.25 0.859375 -0.671875q0.265625 -0.421875 0.265625 -0.9375q0 -0.390625 -0.140625 -0.71875q-0.140625 -0.34375 -0.4375 -0.59375q-0.296875 -0.25 -0.78125 -0.390625q-0.484375 -0.15625 -1.171875 -0.15625l-2.03125 0l0 3.703125zm11.265625 7.015625q-0.5 0 -0.953125 -0.15625q-0.453125 -0.140625 -0.8125 -0.4375q-0.359375 -0.3125 -0.625 -0.78125q-0.265625 -0.46875 -0.40625 -1.125q-0.140625 -0.65625 -0.140625 -1.46875l0.03125 -4.765625l1.296875 0l0 4.765625q0 0.796875 0.140625 1.359375q0.15625 0.5625 0.421875 0.90625q0.28125 0.328125 0.625 0.484375q0.359375 0.140625 0.75 0.140625q0.453125 0 0.859375 -0.171875q0.421875 -0.1875 0.734375 -0.53125q0.328125 -0.359375 0.515625 -0.921875q0.1875 -0.5625 0.1875 -1.328125l0 -4.703125l1.3125 0l0 7.203125q0 0.34375 0.015625 0.671875q0.03125 0.328125 0.171875 0.65625l-1.296875 0q-0.09375 -0.21875 -0.125 -0.4375q-0.03125 -0.21875 -0.046875 -0.4375q-0.015625 -0.21875 -0.015625 -0.4375q-0.203125 0.359375 -0.484375 0.640625q-0.28125 0.28125 -0.625 0.484375q-0.34375 0.1875 -0.734375 0.28125q-0.375 0.109375 -0.796875 0.109375zm8.21875 -7.671875l-1.84375 0l0 -1.0625l1.84375 0l0 -0.703125q0 -1.140625 0.40625 -1.859375q0.40625 -0.734375 1.109375 -1.078125q0.71875 -0.34375 1.625 -0.34375q0.796875 0 1.5 0.296875q0.703125 0.28125 1.171875 0.875l-0.71875 0.984375l-0.109375 0.15625l-0.109375 -0.078125q-0.015625 -0.140625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.3125 -0.265625 -0.609375 -0.375q-0.296875 -0.109375 -0.765625 -0.109375q-0.53125 0 -0.96875 0.21875q-0.421875 0.21875 -0.671875 0.6875q-0.234375 0.46875 -0.234375 1.296875l0 0.703125l2.796875 0l0 1.0625l-2.796875 0l0 7.46875l-1.296875 0l0 -7.46875zm9.328125 0l-1.84375 0l0 -1.0625l1.84375 0l0 -0.703125q0 -1.140625 0.40625 -1.859375q0.40625 -0.734375 1.109375 -1.078125q0.71875 -0.34375 1.625 -0.34375q0.796875 0 1.5 0.296875q0.703125 0.28125 1.171875 0.875l-0.71875 0.984375l-0.109375 0.15625l-0.109375 -0.078125q-0.015625 -0.140625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.3125 -0.265625 -0.609375 -0.375q-0.296875 -0.109375 -0.765625 -0.109375q-0.53125 0 -0.96875 0.21875q-0.421875 0.21875 -0.671875 0.6875q-0.234375 0.46875 -0.234375 1.296875l0 0.703125l2.796875 0l0 1.0625l-2.796875 0l0 7.46875l-1.296875 0l0 -7.46875zm11.390625 7.65625q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625zm9.296875 -3.421875l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875z" fill-rule="nonzero"/><path fill="#0097a7" d="m417.95084 281.98373l261.79526 0l0 37.66928l-261.79526 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.95084 281.98373l261.79526 0l0 37.66928l-261.79526 0z" fill-rule="evenodd"/><path fill="#000000" d="m545.41876 307.9415q-0.453125 0 -0.765625 -0.3125q-0.3125 -0.3125 -0.3125 -0.75q0 -0.421875 0.3125 -0.734375q0.3125 -0.328125 0.765625 -0.328125q0.421875 0 0.734375 0.328125q0.328125 0.3125 0.328125 0.734375q0 0.4375 -0.328125 0.75q-0.3125 0.3125 -0.734375 0.3125zm3.203125 0q-0.4375 0 -0.765625 -0.3125q-0.3125 -0.3125 -0.3125 -0.75q0 -0.421875 0.3125 -0.734375q0.328125 -0.328125 0.765625 -0.328125q0.4375 0 0.75 0.328125q0.3125 0.3125 0.3125 0.734375q0 0.4375 -0.3125 0.75q-0.3125 0.3125 -0.75 0.3125zm3.25 0q-0.453125 0 -0.765625 -0.3125q-0.3125 -0.3125 -0.3125 -0.75q0 -0.421875 0.3125 -0.734375q0.3125 -0.328125 0.765625 -0.328125q0.421875 0 0.734375 0.328125q0.328125 0.3125 0.328125 0.734375q0 0.4375 -0.328125 0.75q-0.3125 0.3125 -0.734375 0.3125z" fill-rule="nonzero"/><path fill="#0097a7" d="m417.95084 319.63815l261.79526 0l0 37.66928l-261.79526 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m417.95084 319.63815l261.79526 0l0 37.66928l-261.79526 0z" fill-rule="evenodd"/><path fill="#000000" d="m521.7547 345.3928l0 -11.625l3.3125 0q1.0 0 1.71875 0.203125q0.734375 0.203125 1.203125 0.578125q0.484375 0.375 0.703125 0.90625q0.234375 0.515625 0.234375 1.125q0 0.546875 -0.203125 1.078125q-0.1875 0.515625 -0.578125 0.90625q-0.375 0.390625 -0.9375 0.59375q0.484375 0.171875 0.875 0.46875q0.40625 0.296875 0.6875 0.703125q0.296875 0.390625 0.453125 0.859375q0.171875 0.453125 0.171875 0.9375q0 0.734375 -0.265625 1.34375q-0.265625 0.59375 -0.796875 1.03125q-0.53125 0.421875 -1.359375 0.65625q-0.8125 0.234375 -1.9375 0.234375l-3.28125 0zm1.296875 -1.140625l2.21875 0q0.8125 0 1.421875 -0.203125q0.625 -0.21875 0.96875 -0.703125q0.34375 -0.484375 0.34375 -1.296875q0 -0.6875 -0.296875 -1.21875q-0.296875 -0.53125 -0.96875 -0.828125q-0.65625 -0.3125 -1.796875 -0.3125l-1.890625 0l0 4.5625zm0 -5.671875l1.828125 0q1.015625 0 1.609375 -0.234375q0.609375 -0.25 0.859375 -0.671875q0.265625 -0.421875 0.265625 -0.9375q0 -0.390625 -0.140625 -0.71875q-0.140625 -0.34375 -0.4375 -0.59375q-0.296875 -0.25 -0.78125 -0.390625q-0.484375 -0.15625 -1.171875 -0.15625l-2.03125 0l0 3.703125zm11.265625 7.015625q-0.5 0 -0.953125 -0.15625q-0.453125 -0.140625 -0.8125 -0.4375q-0.359375 -0.3125 -0.625 -0.78125q-0.265625 -0.46875 -0.40625 -1.125q-0.140625 -0.65625 -0.140625 -1.46875l0.03125 -4.765625l1.296875 0l0 4.765625q0 0.796875 0.140625 1.359375q0.15625 0.5625 0.421875 0.90625q0.28125 0.328125 0.625 0.484375q0.359375 0.140625 0.75 0.140625q0.453125 0 0.859375 -0.171875q0.421875 -0.1875 0.734375 -0.53125q0.328125 -0.359375 0.515625 -0.921875q0.1875 -0.5625 0.1875 -1.328125l0 -4.703125l1.3125 0l0 7.203125q0 0.34375 0.015625 0.671875q0.03125 0.328125 0.171875 0.65625l-1.296875 0q-0.09375 -0.21875 -0.125 -0.4375q-0.03125 -0.21875 -0.046875 -0.4375q-0.015625 -0.21875 -0.015625 -0.4375q-0.203125 0.359375 -0.484375 0.640625q-0.28125 0.28125 -0.625 0.484375q-0.34375 0.1875 -0.734375 0.28125q-0.375 0.109375 -0.796875 0.109375zm8.21875 -7.671875l-1.84375 0l0 -1.0625l1.84375 0l0 -0.703125q0 -1.140625 0.40625 -1.859375q0.40625 -0.734375 1.109375 -1.078125q0.71875 -0.34375 1.625 -0.34375q0.796875 0 1.5 0.296875q0.703125 0.28125 1.171875 0.875l-0.71875 0.984375l-0.109375 0.15625l-0.109375 -0.078125q-0.015625 -0.140625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.3125 -0.265625 -0.609375 -0.375q-0.296875 -0.109375 -0.765625 -0.109375q-0.53125 0 -0.96875 0.21875q-0.421875 0.21875 -0.671875 0.6875q-0.234375 0.46875 -0.234375 1.296875l0 0.703125l2.796875 0l0 1.0625l-2.796875 0l0 7.46875l-1.296875 0l0 -7.46875zm9.328125 0l-1.84375 0l0 -1.0625l1.84375 0l0 -0.703125q0 -1.140625 0.40625 -1.859375q0.40625 -0.734375 1.109375 -1.078125q0.71875 -0.34375 1.625 -0.34375q0.796875 0 1.5 0.296875q0.703125 0.28125 1.171875 0.875l-0.71875 0.984375l-0.109375 0.15625l-0.109375 -0.078125q-0.015625 -0.140625 -0.078125 -0.296875q-0.046875 -0.15625 -0.25 -0.375q-0.3125 -0.265625 -0.609375 -0.375q-0.296875 -0.109375 -0.765625 -0.109375q-0.53125 0 -0.96875 0.21875q-0.421875 0.21875 -0.671875 0.6875q-0.234375 0.46875 -0.234375 1.296875l0 0.703125l2.796875 0l0 1.0625l-2.796875 0l0 7.46875l-1.296875 0l0 -7.46875zm11.390625 7.65625q-0.9375 0 -1.703125 -0.28125q-0.765625 -0.296875 -1.3125 -0.859375q-0.546875 -0.5625 -0.84375 -1.390625q-0.296875 -0.828125 -0.296875 -1.890625q0 -1.078125 0.296875 -1.921875q0.3125 -0.84375 0.84375 -1.40625q0.53125 -0.578125 1.25 -0.875q0.71875 -0.296875 1.515625 -0.296875q0.734375 0 1.359375 0.25q0.640625 0.234375 1.109375 0.734375q0.46875 0.5 0.734375 1.265625q0.265625 0.75 0.265625 1.796875q0 0.15625 -0.015625 0.3125q0 0.15625 -0.015625 0.3125l-6.0625 0q0.03125 0.8125 0.265625 1.421875q0.25 0.609375 0.640625 1.0q0.390625 0.375 0.90625 0.5625q0.515625 0.1875 1.09375 0.1875q0.453125 0 0.84375 -0.09375q0.390625 -0.109375 0.734375 -0.328125q0.34375 -0.21875 0.625 -0.53125l0.75 0.71875q-0.359375 0.4375 -0.8125 0.734375q-0.453125 0.296875 -1.0 0.4375q-0.546875 0.140625 -1.171875 0.140625zm-2.828125 -5.296875l4.6875 0q0 -0.59375 -0.15625 -1.0625q-0.15625 -0.484375 -0.453125 -0.828125q-0.28125 -0.34375 -0.6875 -0.515625q-0.390625 -0.1875 -0.90625 -0.1875q-0.4375 0 -0.84375 0.15625q-0.40625 0.140625 -0.75 0.46875q-0.328125 0.328125 -0.5625 0.8125q-0.234375 0.484375 -0.328125 1.15625zm9.296875 -3.421875l1.40625 0l-0.0625 2.484375l-0.171875 -0.28125q0.109375 -0.578125 0.390625 -1.015625q0.28125 -0.453125 0.671875 -0.75q0.390625 -0.3125 0.859375 -0.46875q0.46875 -0.171875 0.96875 -0.171875q0.671875 0 1.234375 0.234375q0.578125 0.234375 1.03125 0.703125l-0.59375 1.046875l-0.078125 0.171875l-0.125 -0.078125q-0.046875 -0.15625 -0.125 -0.28125q-0.0625 -0.140625 -0.3125 -0.3125q-0.296875 -0.1875 -0.546875 -0.25q-0.25 -0.078125 -0.578125 -0.078125q-0.453125 0 -0.921875 0.1875q-0.46875 0.171875 -0.859375 0.546875q-0.375 0.375 -0.609375 0.953125q-0.21875 0.578125 -0.21875 1.375l0 4.53125l-1.359375 0l0 -8.546875z" fill-rule="nonzero"/><path fill="#000000" fill-opacity="0.0" d="m689.96375 206.21051l0 150.61417" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="3.0" stroke-linejoin="round" stroke-linecap="butt" d="m689.96375 217.58052l0 127.87416" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="3.0" stroke-linecap="butt" d="m685.02875 217.58052l0 -9.869995l9.869995 0l0 9.869995l-9.869995 0z" fill-rule="nonzero"/><path stroke="#ffffff" stroke-width="3.0" stroke-linecap="butt" d="m694.89874 345.45468l0 9.869995l-9.869995 0l0 -9.869995l9.869995 0z" fill-rule="nonzero"/><path fill="#ea9999" d="m549.8757 151.28491l129.88977 0l0 54.92914l-129.88977 0z" fill-rule="evenodd"/><path stroke="#ffffff" stroke-width="1.0" stroke-linejoin="round" stroke-linecap="butt" d="m549.8757 151.28491l129.88977 0l0 54.92914l-129.88977 0z" fill-rule="evenodd"/><path fill="#000000" d="m596.92993 185.66948l0 -11.625l1.03125 0l2.875 5.6875l2.921875 -5.703125l0.984375 0l0 11.640625l-1.234375 0l0 -8.765625l-2.515625 4.6875l-0.5 0l-2.34375 -4.640625l0 8.71875l-1.21875 0zm12.5 0.203125q-1.453125 0 -2.234375 -0.65625q-0.78125 -0.671875 -0.78125 -1.6875q0 -0.6875 0.328125 -1.234375q0.328125 -0.5625 0.921875 -0.9375q0.59375 -0.390625 1.375 -0.59375q0.78125 -0.203125 1.65625 -0.203125q0.3125 0 0.640625 0.015625q0.328125 0 0.6875 0.03125q0.375 0.03125 0.765625 0.078125l0.03125 1.03125q-0.328125 -0.0625 -0.671875 -0.078125q-0.34375 -0.03125 -0.6875 -0.046875q-0.328125 -0.03125 -0.65625 -0.03125q-0.671875 0 -1.234375 0.125q-0.546875 0.109375 -0.96875 0.328125q-0.40625 0.21875 -0.640625 0.578125q-0.21875 0.359375 -0.21875 0.859375q0 0.390625 0.140625 0.671875q0.140625 0.265625 0.375 0.421875q0.25 0.15625 0.578125 0.234375q0.328125 0.0625 0.71875 0.0625q0.734375 0 1.296875 -0.21875q0.578125 -0.234375 0.96875 -0.6875q0.40625 -0.453125 0.609375 -1.09375q0.203125 -0.65625 0.203125 -1.5q0 -1.265625 -0.296875 -1.96875q-0.296875 -0.71875 -0.890625 -1.015625q-0.578125 -0.296875 -1.5 -0.296875q-0.609375 0 -1.203125 0.234375q-0.59375 0.234375 -1.078125 0.71875l-0.65625 -0.859375q0.59375 -0.59375 1.375 -0.890625q0.78125 -0.3125 1.65625 -0.3125q0.84375 0 1.53125 0.21875q0.6875 0.203125 1.1875 0.6875q0.5 0.46875 0.765625 1.265625q0.265625 0.796875 0.265625 2.015625l0 4.53125l-1.453125 0l0 -1.3125q-0.234375 0.46875 -0.59375 0.765625q-0.34375 0.296875 -0.734375 0.46875q-0.390625 0.15625 -0.796875 0.21875q-0.40625 0.0625 -0.78125 0.0625zm10.078125 -0.015625q-0.71875 0 -1.359375 -0.15625q-0.640625 -0.140625 -1.21875 -0.4375q-0.5625 -0.296875 -1.046875 -0.75l0.75 -1.265625l0.109375 -0.171875l0.109375 0.078125q0.046875 0.140625 0.09375 0.3125q0.046875 0.15625 0.25 0.375q0.390625 0.390625 0.9375 0.640625q0.546875 0.25 1.40625 0.25q0.6875 0 1.15625 -0.15625q0.484375 -0.171875 0.75 -0.484375q0.265625 -0.3125 0.265625 -0.71875q0 -0.28125 -0.125 -0.484375q-0.109375 -0.21875 -0.359375 -0.40625q-0.25 -0.1875 -0.671875 -0.375q-0.40625 -0.1875 -1.0 -0.375q-0.96875 -0.28125 -1.703125 -0.65625q-0.71875 -0.375 -1.109375 -0.859375q-0.390625 -0.5 -0.390625 -1.125q0 -0.65625 0.34375 -1.125q0.359375 -0.484375 1.0625 -0.75q0.703125 -0.28125 1.734375 -0.28125q0.59375 0 1.078125 0.09375q0.484375 0.078125 0.890625 0.25q0.40625 0.15625 0.75 0.421875q0.359375 0.25 0.671875 0.578125l-0.765625 0.9375l-0.140625 0.140625l-0.078125 -0.09375q-0.03125 -0.140625 -0.09375 -0.28125q-0.0625 -0.140625 -0.265625 -0.34375q-0.4375 -0.34375 -0.984375 -0.515625q-0.53125 -0.171875 -1.03125 -0.171875q-0.75 0 -1.296875 0.265625q-0.53125 0.25 -0.53125 0.796875q0 0.28125 0.1875 0.546875q0.203125 0.265625 0.71875 0.546875q0.53125 0.265625 1.5 0.546875q1.0625 0.328125 1.703125 0.671875q0.640625 0.34375 0.921875 0.796875q0.296875 0.453125 0.296875 1.0625q0 0.71875 -0.390625 1.328125q-0.375 0.609375 -1.15625 0.984375q-0.78125 0.359375 -1.96875 0.359375l0 0zm8.359375 -4.453125l-1.125 1.03125l0 3.234375l-1.359375 0l0 -12.40625l1.34375 0l0.1875 0l0 0.125q-0.109375 0.125 -0.140625 0.25q-0.03125 0.125 -0.03125 0.453125l0 6.96875l4.265625 -3.953125q0.125 0.015625 0.234375 0.046875q0.109375 0.015625 0.21875 0.03125q0.125 0 0.265625 0.015625q0.140625 0 0.28125 0l0.484375 0l-3.59375 3.375l4.1875 5.09375l-1.734375 0.078125l-3.484375 -4.34375z" fill-rule="nonzero"/></g></svg>
\ No newline at end of file
diff --git a/libc/docs/gpu/rpc.rst b/libc/docs/gpu/rpc.rst
index b02053e2dc72a8..78ae778671881a 100644
--- a/libc/docs/gpu/rpc.rst
+++ b/libc/docs/gpu/rpc.rst
@@ -11,10 +11,298 @@ Remote Procedure Calls
Remote Procedure Call Implementation
====================================
-Certain features from the standard C library, such as allocation or printing,
-require support from the operating system. We instead implement a remote
-procedure call (RPC) interface to allow submitting work from the GPU to a host
-server that forwards it to the host system.
+Traditionally, the C library abstracts over several functions that interface
+with the platform's operating system through system calls. The GPU however does
+not provide an operating system that can handle target dependent operations.
+Instead, we implemented remote procedure calls to interface with the host's
+operating system while executing on a GPU.
+
+We implemented remote procedure calls using unified virtual memory to create a
+shared communicate channel between the two processes. This memory is often
+pinned memory that can be accessed asynchronously and atomically by multiple
+processes simultaneously. This supports means that we can simply provide mutual
+exclusion on a shared better to swap work back and forth between the host system
+and the GPU. We can then use this to create a simple client-server protocol
+using this shared memory.
+
+This work treats the GPU as a client and the host as a server. The client
+initiates a communication while the server listens for them. In order to
+communicate between the host and the device, we simply maintain a buffer of
+memory and two mailboxes. One mailbox is write-only while the other is
+read-only. This exposes three primitive operations: using the buffer, giving
+away ownership, and waiting for ownership. This is implemented as a half-duplex
+transmission channel between the two sides. We decided to assign ownership of
+the buffer to the client when the inbox and outbox bits are equal and to the
+server when they are not.
+
+In order to make this transmission channel thread-safe, we abstract ownership of
+the given mailbox pair and buffer around a port, effectively acting as a lock
+and an index into the allocated buffer slice. The server and device have
+independent locks around the given port. In this scheme, the buffer can be used
+to communicate intent and data generically with the server. We them simply
+provide multiple copies of this protocol and expose them as multiple ports.
+
+If this were simply a standard CPU system, this would be sufficient. However,
+GPUs have my unique architectural challenges. First, GPU threads execute in
+lock-step with each other in groups typically called warps or wavefronts. We
+need to target the smallest unit of independent parallelism, so the RPC
+interface needs to handle an entire group of threads at once. This is done by
+increasing the size of the buffer and adding a thread mask argument so the
+server knows which threads are active when it handles the communication. Second,
+GPUs generally have no forward progress guarantees. In order to guarantee we do
+not encounter deadlocks while executing it is required that the number of ports
+matches the maximum amount of hardware parallelism on the device. It is also
+very important that the thread mask remains consistent while interfacing with
+the port.
+
+.. image:: ./rpc-diagram.svg
+ :width: 75%
+ :align: center
+
+The above diagram outlines the architecture of the RPC interface. For clarity
+the following list will explain the operations done by the client and server
+respectively when initiating a communication.
+
+First, a communication from the perspective of the client:
+
+* The client searches for an available port and claims the lock.
+* The client checks that the port is still available to the current device and
+ continues if so.
+* The client writes its data to the fixed-size packet and toggles its outbox.
+* The client waits until its inbox matches its outbox.
+* The client reads the data from the fixed-size packet.
+* The client closes the port and continues executing.
+
+Now, the same communication from the perspective of the server:
+
+* The server searches for an available port with pending work and claims the
+ lock.
+* The server checks that the port is still available to the current device.
+* The server reads the opcode to perform the expected operation, in this
+ case a receive and then send.
+* The server reads the data from the fixed-size packet.
+* The server writes its data to the fixed-size packet and toggles its outbox.
+* The server closes the port and continues searching for ports that need to be
+ serviced
+
+This architecture currently requires that the host periodically checks the RPC
+server's buffer for ports with pending work. Note that a port can be closed
+without waiting for its submitted work to be completed. This allows us to model
+asynchronous operations that do not need to wait until the server has completed
+them. If an operation requires more data than the fixed size buffer, we simply
+send multiple packets back and forth in a streaming fashion.
+
+Server Library
+--------------
+
+The RPC server's basic functionality is provided by the LLVM C library. A static
+library called ``libllvmlibc_rpc_server.a`` includes handling for the basic
+operations, such as printing or exiting. This has a small API that handles
+setting up the unified buffer and an interface to check the opcodes.
+
+Some operations are too divergent to provide generic implementations for, such
+as allocating device accessible memory. For these cases, we provide a callback
+registration scheme to add a custom handler for any given opcode through the
+port API. More information can be found in the installed header
+``<install>/include/gpu-none-llvm/rpc_server.h``.
+
+Client Example
+--------------
+
+The Client API is not currently exported by the LLVM C library. This is
+primarily due to being written in C++ and relying on internal data structures.
+It uses a simple send and receive interface with a fixed-size packet. The
+following example uses the RPC interface to call a function pointer on the
+server.
+
+This code first opens a port with the given opcode to facilitate the
+communication. It then copies over the argument struct to the server using the
+``send_n`` interface to stream arbitrary bytes. The next send operation provides
+the server with the function pointer that will be executed. The final receive
+operation is a no-op and simply forces the client to wait until the server is
+done. It can be omitted if asynchronous execution is desired.
+
+.. code-block:: c++
+
+ void rpc_host_call(void *fn, void *data, size_t size) {
+ rpc::Client::Port port = rpc::client.open<RPC_HOST_CALL>();
+ port.send_n(data, size);
+ port.send([=](rpc::Buffer *buffer) {
+ buffer->data[0] = reinterpret_cast<uintptr_t>(fn);
+ });
+ port.recv([](rpc::Buffer *) {});
+ port.close();
+ }
+
+Server Example
+--------------
+
+This example shows the server-side handling of the previous client example. When
+the server is checked, if there are any ports with pending work it will check
+the opcode and perform the appropriate action. In this case, the action is to
+call a function pointer provided by the client.
+
+In this example, the server simply runs forever in a separate thread for
+brevity's sake. Because the client is a GPU potentially handling several threads
+at once, the server needs to loop over all the active threads on the GPU. We
+abstract this into the ``lane_size`` variable, which is simply the device's warp
+or wavefront size. The identifier is simply the threads index into the current
+warp or wavefront. We allocate memory to copy the struct data into, and then
+call the given function pointer with that copied data. The final send simply
+signals completion and uses the implicit thread mask to delete the temporary
+data.
+
+.. code-block:: c++
+
+ for(;;) {
+ auto port = server.try_open(index);
+ if (!port)
+ return continue;
+
+ switch(port->get_opcode()) {
+ case RPC_HOST_CALL: {
+ uint64_t sizes[LANE_SIZE];
+ void *args[LANE_SIZE];
+ port->recv_n(args, sizes, [&](uint64_t size) { return new char[size]; });
+ port->recv([&](rpc::Buffer *buffer, uint32_t id) {
+ reinterpret_cast<void (*)(void *)>(buffer->data[0])(args[id]);
+ });
+ port->send([&](rpc::Buffer *, uint32_t id) {
+ delete[] reinterpret_cast<uint8_t *>(args[id]);
+ });
+ break;
+ }
+ default:
+ port->recv([](rpc::Buffer *) {});
+ break;
+ }
+ }
+
+CUDA Server Example
+-------------------
+
+The following code shows an example of using the exported RPC interface along
+with the C library to manually configure a working server using the CUDA
+language. Other runtimes can use the presence of the ``__llvm_libc_rpc_client``
+in the GPU executable as an indicator for whether or not the server can be
+checked. These details should ideally be handled by the GPU language runtime,
+but the following example shows how it can be used by a standard user.
+
+.. code-block:: cuda
+
+ #include <cstdio>
+ #include <cstdlib>
+ #include <cuda_runtime.h>
+
+ #include <gpu-none-llvm/rpc_server.h>
+
+ [[noreturn]] void handle_error(cudaError_t err) {
+ fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
+ exit(EXIT_FAILURE);
+ }
+
+ [[noreturn]] void handle_error(rpc_status_t err) {
+ fprintf(stderr, "RPC error: %d\n", err);
+ exit(EXIT_FAILURE);
+ }
+
+ // The handle to the RPC client provided by the C library.
+ extern "C" __device__ void *__llvm_libc_rpc_client;
+
+ __global__ void get_client_ptr(void **ptr) { *ptr = __llvm_libc_rpc_client; }
+
+ // Obtain the RPC client's handle from the device. The CUDA language cannot look
+ // up the symbol directly like the driver API, so we launch a kernel to read it.
+ void *get_rpc_client() {
+ void *rpc_client = nullptr;
+ void **rpc_client_d = nullptr;
+
+ if (cudaError_t err = cudaMalloc(&rpc_client_d, sizeof(void *)))
+ handle_error(err);
+ get_client_ptr<<<1, 1>>>(rpc_client_d);
+ if (cudaError_t err = cudaDeviceSynchronize())
+ handle_error(err);
+ if (cudaError_t err = cudaMemcpy(&rpc_client, rpc_client_d, sizeof(void *),
+ cudaMemcpyDeviceToHost))
+ handle_error(err);
+ return rpc_client;
+ }
+
+ // Routines to allocate mapped memory that both the host and the device can
+ // access asychonrously to communicate with eachother.
+ void *alloc_host(size_t size, void *) {
+ void *sharable_ptr;
+ if (cudaError_t err = cudaMallocHost(&sharable_ptr, sizeof(void *)))
+ handle_error(err);
+ return sharable_ptr;
+ };
+
+ void free_host(void *ptr, void *) {
+ if (cudaError_t err = cudaFreeHost(ptr))
+ handle_error(err);
+ }
+
+ // The device-side overload of the standard C function to call.
+ extern "C" __device__ int puts(const char *);
+
+ // Calls the C library function from the GPU C library.
+ __global__ void hello() { puts("Hello world!"); }
+
+ int main() {
+ int device = 0;
+ // Initialize the RPC server to run on a single device.
+ if (rpc_status_t err = rpc_init(/*num_device=*/1))
+ handle_error(err);
+
+ // Initialize the RPC server to run on the given device.
+ if (rpc_status_t err =
+ rpc_server_init(device, RPC_MAXIMUM_PORT_COUNT,
+ /*warp_size=*/32, alloc_host, /*data=*/nullptr))
+ handle_error(err);
+
+ // Initialize the RPC client by copying the buffer to the device's handle.
+ void *rpc_client = get_rpc_client();
+ if (cudaError_t err =
+ cudaMemcpy(rpc_client, rpc_get_client_buffer(device),
+ rpc_get_client_size(), cudaMemcpyHostToDevice))
+ handle_error(err);
+
+ cudaStream_t stream;
+ if (cudaError_t err = cudaStreamCreate(&stream))
+ handle_error(err);
+
+ // Execute the kernel.
+ hello<<<1, 1, 0, stream>>>();
+
+ // While the kernel is executing, check the RPC server for work to do.
+ while (cudaStreamQuery(stream) == cudaErrorNotReady)
+ if (rpc_status_t err = rpc_handle_server(device))
+ handle_error(err);
+
+ // Shut down the server running on the given device.
+ if (rpc_status_t err =
+ rpc_server_shutdown(device, free_host, /*data=*/nullptr))
+ handle_error(err);
+
+ // Shut down the entire RPC server interface.
+ if (rpc_status_t err = rpc_shutdown())
+ handle_error(err);
+
+ return EXIT_SUCCESS;
+ }
+
+The above code must be compiled in CUDA's relocatable device code mode and with
+the advanced offloading driver to link in the library. Currently this can be
+done with the following invocation. Using LTO avoids the overhead normally
+associated with relocatable device code linking.
+
+.. code-block:: sh
+
+ $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart -lcgpu \
+ -I<install-path>include -L<install-path>/lib -lllvmlibc_rpc_server \
+ -O3 -foffload-lto -o hello
+ $> ./hello
+ Hello world!
Extensions
----------
diff --git a/libc/docs/gpu/testing.rst b/libc/docs/gpu/testing.rst
index d0f162694562de..f793e938e9806f 100644
--- a/libc/docs/gpu/testing.rst
+++ b/libc/docs/gpu/testing.rst
@@ -18,9 +18,9 @@ Testing Infrastructure
======================
The testing support in LLVM's libc implementation for GPUs is designed to mimic
-the standard unit tests as much as possible. We use the `remote procedure call
-<libc_gpu_rpc>`_ support to provide the necessary utilities like printing from
-the GPU. Execution is performed by emitting a ``_start`` kernel from the GPU
+the standard unit tests as much as possible. We use the :ref:`libc_gpu_rpc`
+support to provide the necessary utilities like printing from the GPU. Execution
+is performed by emitting a ``_start`` kernel from the GPU
that is then called by an external loader utility. This is an example of how
this can be done manually:
More information about the libc-commits
mailing list