Remote Compilation

A GPU server can compile and evaluate a kernel in one request, as shown in Benchmark a Kernel with KCoral. When compilation takes much longer than execution, however, moving builds to CPU-only machines lets you scale compilation capacity independently of your GPU machines.

In this tutorial, you will compile a CUDA C kernel on a CPU server and execute it on a separate GPU server. The first request returns the compiled shared library as bytes to the client. The client then uploads those bytes to the GPU server in a second request, which checks the kernel’s output and measures its execution time.

This workflow requires a compiler that can build without accessing a GPU. Compiler APIs that query CUDA or load GPU modules during compilation still need a GPU server.

The client carries the compiled artifact between the requests. Each request has its own Program:

Request	Destination	Inputs	Returned values
1. Compile	CPU server	Kernel source and the execution GPU’s target architecture	Shared-library bytes
2. Execute	GPU server	Those library bytes and the workload	Correctness and timing reports

You will follow this two-request workflow with a client that compiles a CUDA C add-one kernel and then checks and benchmarks it.

Prepare the two servers

The CPU server needs build tools, while the GPU server needs the GPU runtime and measurement tools. Prepare each environment for its role. Install the client on the machine coordinating the requests. It needs no local compiler or GPU. The two server roles have different requirements:

Role	Environment
CPU compiler	The compiler environment, the CUDA toolkit with nvcc, and a host C++ compiler
GPU executor	The GPU worker environment, including PyTorch, TVM FFI and CUPTI for benchmarking

The compiled library must be compatible with the GPU server’s operating system, CPU architecture, GPU architecture, and runtime dependencies. Keep the two servers’ TVM FFI and CUDA components compatible. See the library upload protocol for loading and export requirements.

This example uploads Python that builds CUDA C through TVM FFI, then uses CUPTI to collect GPU activity timestamps.

Warning

KCoral allows clients to execute arbitrary code on its workers. Only allow trusted clients to access your KCoral server or Router. Deploy on a trusted, isolated network and never expose these endpoints to the public internet. Run workers in a sandbox with restricted permissions and access to host resources.

For a local demonstration, start a CPU server in one terminal:

kcoral server --device cpu --num-workers 8 --host 127.0.0.1 --port 8000

In another terminal on the same machine, start a GPU server:

kcoral server --device gpu --gpus 0 --host 127.0.0.1 --port 8001

Both commands listen only on the local machine. To run the servers on separate machines, use --host 0.0.0.0 on each server to accept connections over a trusted network, and use their reachable addresses in the client command below. See Launch the server for configuration.

Run the complete client

First run the whole workflow to see the two servers working together. The sections that follow walk through how the client constructs each request. From the repository checkout, run the client with both servers on the local machine, or replace 127.0.0.1 with each server’s reachable address:

KCORAL_CPU_URL=http://127.0.0.1:8000 \
KCORAL_GPU_URL=http://127.0.0.1:8001 \
python examples/cpu_compile/cpu_compile_gpu_execute.py

Download the complete client. The downloaded file can be run directly with the same environment variables. If both servers run on the client machine, the defaults are http://127.0.0.1:8000 for compilation and http://127.0.0.1:8001 for execution.

On success, the script prints three lines with this format; the size, architecture, and GPU timings depend on your environment:

compiled <size> KiB for <arch>
CPU request held a GPU lease for 0 ms
GPU request held its lease for <time> ms; kernel median <latency> us

The CPU request compiles without using a GPU. The GPU request checks correctness before measuring the kernel, so reaching this output means both requests completed and the correctness check passed. A GPU lease is exclusive access to the device; the GPU request’s lease time covers more than the kernel measurement. If either program returns FAILED, the script instead reports compile failed: ... or benchmark failed: ... with the error details. HTTP and connection failures raise client exceptions.

Now that you have seen the complete workflow, the following sections trace how the client chooses a compilation target, builds the library, and submits it to the GPU server.

Read the GPU target

The compiler needs to know what GPU architecture to build for. The CPU server has no GPU from which to determine that architecture, so the client asks the GPU server before building:

arch = gpu_client.target()["arch"]

The client passes arch to the uploaded compile_cuda_binary function so the library is built for the GPU that will execute it.

Compile and return the library

The first request turns source code into a library that the client can carry to the GPU server. The source defines a GPU kernel and a host function, add_one, which launches it. The host function uses TVM FFI’s TensorView parameters to access tensors. The uploaded compiler uses tvm_ffi.cpp.build_inline to supply the required includes and export wrapper.

CUDA_SOURCE = r"""
__global__ void add_one_kernel(const float* x, float* y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = x[i] + 1.0f;
}

void add_one(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {
  int n = static_cast<int>(x.numel());
  add_one_kernel<<<(n + 255) / 256, 256>>>(static_cast<const float*>(x.data_ptr()),
                                           static_cast<float*>(y.data_ptr()), n);
}
"""

The example’s OPERATIONS string defines the Python compiler and execution helpers uploaded by each program:

OPERATIONS = r"""
from kcoral.builtins import benchmark


def empty(spec):
    import torch

    return torch.empty(spec["shape"], dtype=getattr(torch, spec["dtype"]), device="cuda")


def randn(spec):
    import torch

    generator = torch.Generator(device="cuda").manual_seed(spec.get("seed", 0))
    return torch.randn(
        spec["shape"], dtype=getattr(torch, spec["dtype"]), device="cuda", generator=generator
    )


def assert_close(actual, expected):
    import torch

    torch.testing.assert_close(actual.cpu(), expected.cpu(), rtol=1e-2, atol=1e-3)
    return {"ok": True}

Show full sourceShow less: cpu_compile/cpu_compile_gpu_execute.py

OPERATIONS = r"""
from kcoral.builtins import benchmark


def empty(spec):
    import torch

    return torch.empty(spec["shape"], dtype=getattr(torch, spec["dtype"]), device="cuda")


def randn(spec):
    import torch

    generator = torch.Generator(device="cuda").manual_seed(spec.get("seed", 0))
    return torch.randn(
        spec["shape"], dtype=getattr(torch, spec["dtype"]), device="cuda", generator=generator
    )


def assert_close(actual, expected):
    import torch

    torch.testing.assert_close(actual.cpu(), expected.cpu(), rtol=1e-2, atol=1e-3)
    return {"ok": True}


def compile_cuda_binary(source_path, cfg):
    import os
    from pathlib import Path

    import tvm_ffi.cpp

    arch = cfg["arch"].removeprefix("sm_")
    suffix = "a" if arch.endswith("a") else ""
    digits = arch.removesuffix("a")
    key = "TVM_FFI_CUDA_ARCH_LIST"
    previous = os.environ.get(key)
    os.environ[key] = f"{int(digits[:-1])}.{digits[-1]}{suffix}"
    try:
        path = tvm_ffi.cpp.build_inline(
            name=cfg["name"],
            cuda_sources=Path(source_path).read_text(encoding="utf-8"),
            functions=cfg["functions"],
            backend="cuda",
            extra_cuda_cflags=cfg.get("extra_cuda_cflags"),
        )
        return Path(path).read_bytes()
    finally:
        if previous is None:
            os.environ.pop(key, None)
        else:
            os.environ[key] = previous
"""

The compilation program uploads that Python and the CUDA source file, selects compile_cuda_binary with get_function(..., cpu_only=True), passes the file path and configuration (functions=["add_one"], arch), and returns the library:

def compile_program(arch: str) -> Program:
    program = Program()
    operations = program.upload(kind="module", source=OPERATIONS)
    compile_cuda_binary = program.get_function(
        module=operations, name="compile_cuda_binary", cpu_only=True
    )
    source_path = program.upload_file(blob=CUDA_SOURCE.encode("utf-8"), path="src/add_one.cu")
    library = program.run(
        fn=compile_cuda_binary,
        args=[
            source_path,
            {
                "name": "example_add_one",
                "functions": ["add_one"],
                "arch": arch,
                "extra_cuda_cflags": ["-O3"],
            },
        ],
    )
    program.return_(key="library", value=library)
    return program

The uploaded compile_cuda_binary function builds a shared library without loading it or launching the kernel. return_(key="library", ...) selects the bytes for the response. After cpu_client.execute() succeeds, compiled.results["library"] is a Python bytes object containing the shared library, equivalent to the contents of a .so file.

Upload and benchmark the library

With compilation complete, the second request can focus on checking and measuring the kernel. Its program receives those bytes as its library argument. It uploads them with kind="library", selects the exported function, creates input and output tensors, and runs the kernel:

def benchmark_program(library: bytes) -> Program:
    program = Program()
    operations = program.upload(kind="module", source=OPERATIONS)
    empty = program.get_function(module=operations, name="empty")
    randn = program.get_function(module=operations, name="randn")
    assert_close = program.get_function(module=operations, name="assert_close")
    benchmark = program.get_function(module=operations, name="benchmark")
    module = program.upload(kind="library", value=library)
    kernel = program.get_function(module=module, name="add_one")
    reference_module = program.upload(kind="module", source=REFERENCE)
    reference = program.get_function(module=reference_module, name="main")
    src = program.run(
        fn=randn,
        args=[{"shape": [N], "dtype": "float32", "seed": 0}],
    )
    dst = program.run(fn=empty, args=[{"shape": [N], "dtype": "float32"}])
    program.run(fn=kernel, args=[src, dst])
    expected = program.run(fn=reference, args=[src])
    check = program.run(fn=assert_close, args=[dst, expected])
    timing = program.run(
        fn=benchmark,
        args=[kernel, src, dst, {"warmup_ms": 25, "repeat_ms": 100}],
    )
    program.return_(key="check", value=check)
    program.return_(key="timing", value=timing)
    return program

The GPU server loads the uploaded library, and get_function selects its exported add_one function. Calling it with program.run launches the kernel that was compiled on the CPU server. The program then checks the output and benchmarks the kernel only if that check passes, following the benchmark tutorial.

Submit both programs

The client’s main function brings these steps together. It connects to both servers, runs the compilation program, and uses the returned library to run the benchmark program:

def main() -> None:
    cpu_url = os.environ.get("KCORAL_CPU_URL", "http://127.0.0.1:8000")
    gpu_url = os.environ.get("KCORAL_GPU_URL", "http://127.0.0.1:8001")
    with Client(cpu_url) as cpu_client, Client(gpu_url) as gpu_client:
        arch = gpu_client.target()["arch"]
        compiled = cpu_client.execute(compile_program(arch), timeout_seconds=120)
        if not compiled.completed:
            raise SystemExit(f"compile failed: {compiled.error}")
        library = compiled.results["library"]
        if not isinstance(library, bytes):
            raise SystemExit("compile server returned a non-binary library")

        result = gpu_client.execute(benchmark_program(library), timeout_seconds=120)
        if not result.completed:
            raise SystemExit(f"benchmark failed: {result.error}")

    timing = result.results["timing"]
    print(f"compiled {len(library) / 1024:.0f} KiB for {arch}")
    print(f"CPU request held a GPU lease for {compiled.lease_held_ms:.0f} ms")
    print(
        f"GPU request held its lease for {result.lease_held_ms:.0f} ms; "
        f"kernel median {timing['latency_ms_median'] * 1e3:.1f} us"
    )

The client checks each result before proceeding: a failed compilation stops it before the GPU submission, and a failed GPU program stops it before printing a successful timing report. Inspect the result’s error and request_id to identify the failing instruction and find the corresponding server logs. See Read results for handling program failures and connection exceptions.

After the GPU server finishes, the client reads the timing report from result.results["timing"] and prints the kernel’s median duration. The correctness report is also available in result.results["check"]. For the meaning of the timing statistics, see Measure GPU activity.