[all-commits] [llvm/llvm-project] f879ac: [libc] Rework the RPC interface to accept runtime ...

[Joseph Huber via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: f879ac0385d4c5f7b2b9f4807cd7bd4a78556c1c
      https://github.com/llvm/llvm-project/commit/f879ac0385d4c5f7b2b9f4807cd7bd4a78556c1c
  Author: Joseph Huber <huberjn at outlook.com>
  Date:   2024-02-13 (Tue, 13 Feb 2024)

  Changed paths:
    M libc/src/__support/GPU/amdgpu/utils.h
    M libc/src/__support/GPU/generic/utils.h
    M libc/src/__support/GPU/nvptx/utils.h
    M libc/src/__support/RPC/rpc.h
    M libc/test/src/__support/RPC/rpc_smoke_test.cpp
    M libc/utils/gpu/server/rpc_server.cpp

  Log Message:
  -----------
  [libc] Rework the RPC interface to accept runtime wave sizes (#80914)

Summary:
The RPC interface needs to handle an entire warp or wavefront at once.
This is currently done by using a compile time constant indicating the
size of the buffer, which right now defaults to some value on the client
(GPU) side. However, there are currently attempts to move the `libc`
library to a single IR build. This is problematic as the size of the
wave fronts changes between ISAs on AMDGPU. The builitin
`__builtin_amdgcn_wavefrontsize()` will return the appropriate value,
but it is only known at runtime now.

In order to support this, this patch restructures the packet. Now
instead of having an array of arrays, we simply have a large array of
buffers and slice it according to the runtime value if we don't know it
ahead of time. This also somewhat has the advantage of making the buffer
contiguous within a page now that the header has been moved out of it.