[llvm] [Offload] Move RPC server handling to a dedicated thread (PR #112988)

[Joseph Huber via llvm-commits]

================
@@ -1294,10 +1275,26 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
                     reinterpret_cast<void *>(&LaunchParams.Size),
                     CU_LAUNCH_PARAM_END};
 
+  // If we are running an RPC server we want to wake up the server thread
+  // whenever there is a kernel running and let it sleep otherwise.
+  if (GenericDevice.getRPCServer())
+    GenericDevice.Plugin.getRPCServer().Thread->notify();
+
   CUresult Res = cuLaunchKernel(Func, NumBlocks, /*gridDimY=*/1,
                                 /*gridDimZ=*/1, NumThreads,
                                 /*blockDimY=*/1, /*blockDimZ=*/1,
                                 MaxDynCGroupMem, Stream, nullptr, Config);
+
+  // Register a callback to indicate when the kernel is complete.
+  if (GenericDevice.getRPCServer())
+    cuStreamAddCallback(
----------------
jhuber6 wrote:

Didn't know about that, supposedly they have `cuLaunchHostFunc` which just pushes the function onto the stream. It appears the only difference is that the function will not be handled on stream error. I wonder what the proper way to handle the errors are then.

https://github.com/llvm/llvm-project/pull/112988