[Renaud-K]

Let an application bring its own GPU dispatch kernel

Summary

The dispatcher's built-in unified kernel runs one block, and every thread of
that block calls each handler. That suits most handlers, but not every
application: some want a single thread, a cooperative grid, or clusters.

This PR lets an application supply its own dispatch kernel. It reuses the
run-time link introduced in #5722, so the custom kernel ends up in the same
CUDA module as the bridge's hooks and the application's handlers. From the
application's side, it takes three steps: write the kernel, add one CMake
line, and make two calls at run time.

Bringing your own dispatcher in three steps

1. Write the kernel. It is a plain .cu file with one
extern "C" __global__ function. One thread drives the transport through the
same three hooks the built-in kernel uses (cudaq_dev_enter_kernel,
cudaq_dev_rx_poll, cudaq_dev_tx_publish), and the kernel decides how the
handlers run. The framing helpers the built-in kernel uses
(unified_lookup_entry, write_response_header) are shared in
unified_dispatch_device.cuh.

2. Build it with one CMake call, next to the binary that holds the
handlers:

cudaq_realtime_add_dispatch_kernel(my_app my_dispatcher.cu)

The function comes with the cudaq-realtime package config. The kernel's file
is not listed among the target's sources.

3. Look it up and hand the dispatcher a launcher, after
cudaq_dispatcher_set_bridge:

cudaKernel_t kernel;
cudaq_dispatcher_get_device_kernel(dispatcher, "my_dispatcher", &kernel);
cudaq_dispatcher_set_unified_launch(dispatcher, my_launcher, /*ctx=*/NULL);
cudaq_dispatcher_start(dispatcher);

cudaq_dispatcher_start then calls my_launcher instead of launching the
built-in kernel. The launcher receives everything the kernel needs (the
bridge's transport context, the slot size, the function table, the shutdown
flag, the stats counter and the stream) and launches the kernel any way it
likes: any grid, any launch attributes. It launches on the given stream and
returns CUDAQ_OK.

That is all. The bridge, the handlers and the rest of the wiring don't change.

Kernels and handlers: same lookup, different results

cudaq_dispatcher_get_device_kernel is the twin of
cudaq_dispatcher_get_device_handler:

• The same: both find code from the application's own binary in the
  dispatcher's linked module. Whichever lookup comes first performs the link,
  and the handlers and the kernel of one dispatcher must come from the same
  binary.
• The difference: a handler lookup returns a device address, which goes
  into the function table and is called from device code. CUDA has no way to
  find a device function by name, so the handler must be exported with
  CUDAQ_REALTIME_EXPORT_DEVICE_FUNCTION. A kernel lookup returns a host-side
  handle (cudaKernel_t) that the host uses to launch the kernel. A kernel is
  found by its own name, so it needs no export, and it never goes in the
  function table.

Example: the test server

cudaq-realtime-test-server --path=device --dispatch=unified --dispatcher=single_thread_dispatcher
runs the whole flow, and is what the new two-process test shape uses.

• The kernel, unittests/dispatchers/single_thread_dispatcher.cu, is about
  40 lines and does the simplest possible thing. A single thread opens the
  transport session, then loops: wait for a request, look up its handler in
  the function table, call it, write the response header, publish. There are
  no block barriers and no shared-memory hand-off: the whole dispatch loop is
  visible at a glance.
• The build adds it with one line in
  unittests/test_server/CMakeLists.txt:
  cudaq_realtime_add_dispatch_kernel(cudaq-realtime-test-server .../single_thread_dispatcher.cu).
• The wiring is step [5k] of the server: one call looks up the kernel,
  and one call installs the launcher. The launcher, launch_own_dispatcher,
  passes the kernel's arguments in order and launches one block of one thread
  through the driver API (cuLaunchKernelEx), which takes the handle on both
  CUDA 12 and 13.

Under the hood: from a .cu file to a launchable kernel

flowchart TB
  subgraph build ["Build time"]
    direction LR
    src["my_dispatcher.cu"] -->|"nvcc -fatbin -rdc, code=lto_80"| fatbin["LTO-IR fatbin"]
    fatbin -->|"asm .incbin"| section["app binary: __nv_relfatbin section, next to the handlers"]
  end
  subgraph runtime ["Run time"]
    direction LR
    link["nvJitLink -lto: built-in kernel + bridge hooks + app code"] --> lib["CUlibrary: one CUDA module"]
    lib -->|"cudaq_dispatcher_get_device_kernel: cuLibraryGetKernel"| handle["cudaKernel_t"]
    handle -->|"start: the app's launcher"| run["kernel running"]
  end
  build -->|"first lookup: cudaq_device_code_of_this_binary"| runtime

Loading

Why it can't be an ordinary source. The kernel calls the bridge's hooks,
which only the bridge defines. If the kernel were a normal source of the
application, the application's own device link would fail on those undefined
hooks. Weak stubs don't help either: nvcc ignores __attribute__((weak)) on
device functions, so a stub would be linked in place of the bridge's hook.

What cudaq_realtime_add_dispatch_kernel does instead.

• It compiles the .cu file on its own into an LTO-IR fatbin, with the
  target's include directories, compile definitions and C++ standard. By
  default it targets the architecture of the built-in kernel; LTO-IR built for
  one architecture links for that GPU and every newer one.
• It generates a small .cpp file that uses .incbin to place the fatbin in
  the binary's __nv_relfatbin section, next to the fatbins nvcc emits for the
  handlers. No host code registers it, so the binary's own device link never
  sees it.

At run time, the first lookup reads the binary's whole __nv_relfatbin
section with cudaq_device_code_of_this_binary. That section now holds the
handlers and the custom kernel. As in #5722, the dispatcher links that code
with its built-in kernel and the bridge's hooks into one module, with
nvJitLink for the GPU it runs on, and loads it. Because everything is LTO-IR,
the hooks are inlined into the custom kernel just as they are into the
built-in one. cudaq_dispatcher_get_device_kernel then fetches the custom
kernel from that module by name (cuLibraryGetKernel), and start hands
control to the application's launcher. start still checks that every
DEVICE_CALL entry holds an address from the linked module, but no longer
checks config.num_blocks: the launcher chooses the grid.

API changes

• New:
  • cudaq_dispatcher_get_device_kernel (header-only, like
    cudaq_dispatcher_get_device_handler) and the exported
    cudaq_dispatcher_get_device_kernel_in.
  • The CMake function cudaq_realtime_add_dispatch_kernel, installed with
    the package config.
  • unified_dispatch_device.cuh, with the framing helpers shared by the
    built-in kernel and custom ones.
• Changed:
  • cudaq_dispatcher_set_unified_launch also works over a bridge's GPU data
    plane. It takes a NULL transport context there, because the bridge's
    context is kept.
  • cudaq_unified_launch_fn_t now returns a cudaq_status_t, and a failure
    fails cudaq_dispatcher_start. A bridge launch override written against
    the old void form must add the return value.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[Renaud-K]

/ok to test fc401d8

Command Bot: Processing...

[github-actions]

CI Summary (push) — ✅ passed

Run #40982538492 · ✅ 8 · ⏩ 10 · ❌ 0 · ⛔ 0

Top-level jobs (18)

Job	Result
binaries	⏩ skipped
build_and_test	✅ success
changes	✅ success
config_devdeps	✅ success
config_source_build	⏩ skipped
config_wheeldeps	✅ success
devdeps	✅ success
docker_image	⏩ skipped
documentation	✅ success
gen_code_coverage	⏩ skipped
metadata	✅ success
python_devel_wheel	⏩ skipped
python_metapackages	⏩ skipped
python_wheels	⏩ skipped
python_wheels_cu12_validation	⏩ skipped
python_wheels_cu13	⏩ skipped
source_build	⏩ skipped
wheeldeps	✅ success

⏩ Skipped jobs (10) — intentionally skipped on PR builds; run on merge_group / workflow_dispatch

Job
binaries
config_source_build
docker_image
gen_code_coverage
python_devel_wheel
python_metapackages
python_wheels_cu12_validation
python_wheels_cu13
python_wheels
source_build

All sub-jobs (47) — every matrix leg, with links

Job	Status	Link
Build and test (amd64, gcc12, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (amd64, gcc12, openmpi) / Dev environment (Python)	✅ success	view
Build and test (amd64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (amd64, llvm, openmpi) / Dev environment (Python)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Python)	✅ success	view
CI Summary	❔ in_progress	view
Check for stable CUDA-Q changes	✅ success	view
Configure build (devdeps)	✅ success	view
Configure build (source_build)	⏩ skipped	view
Configure build (wheeldeps)	✅ success	view
Create CUDA 12 Python wheels	⏩ skipped	view
Create CUDA 13 Python wheels	⏩ skipped	view
Create CUDA Quantum installer	⏩ skipped	view
Create Docker images	⏩ skipped	view
Create Python metapackages	⏩ skipped	view
Create shared Python development wheel	⏩ skipped	view
Documentation / Build and check docs	✅ success	view
Gen code coverage	⏩ skipped	view
Load dependencies (amd64, gcc12) / Caching	✅ success	view
Load dependencies (amd64, gcc12) / Finalize	✅ success	view
Load dependencies (amd64, gcc12) / Metadata	✅ success	view
Load dependencies (amd64, llvm) / Caching	✅ success	view
Load dependencies (amd64, llvm) / Finalize	✅ success	view
Load dependencies (amd64, llvm) / Metadata	✅ success	view
Load dependencies (arm64, gcc12) / Caching	✅ success	view
Load dependencies (arm64, gcc12) / Finalize	✅ success	view
Load dependencies (arm64, gcc12) / Metadata	✅ success	view
Load dependencies (arm64, llvm) / Caching	✅ success	view
Load dependencies (arm64, llvm) / Finalize	✅ success	view
Load dependencies (arm64, llvm) / Metadata	✅ success	view
Load source build cache	⏩ skipped	view
Load wheel dependencies (amd64, 12.6) / Caching	✅ success	view
Load wheel dependencies (amd64, 12.6) / Finalize	✅ success	view
Load wheel dependencies (amd64, 12.6) / Metadata	✅ success	view
Load wheel dependencies (amd64, 13.0) / Caching	✅ success	view
Load wheel dependencies (amd64, 13.0) / Finalize	✅ success	view
Load wheel dependencies (amd64, 13.0) / Metadata	✅ success	view
Load wheel dependencies (arm64, 12.6) / Caching	✅ success	view
Load wheel dependencies (arm64, 12.6) / Finalize	✅ success	view
Load wheel dependencies (arm64, 12.6) / Metadata	✅ success	view
Load wheel dependencies (arm64, 13.0) / Caching	✅ success	view
Load wheel dependencies (arm64, 13.0) / Finalize	✅ success	view
Load wheel dependencies (arm64, 13.0) / Metadata	✅ success	view
Prepare cache clean-up	❔ in_progress	view
Retrieve PR info	✅ success	view
Validate CUDA 12 Python wheels	⏩ skipped	view

✅ Required checks (7/7) — declared in .github/required-checks.yml for push

Required check	Status	Link
Build and test (amd64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (amd64, llvm, openmpi) / Dev environment (Python)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Python)	✅ success	view
Build and test (amd64, gcc12, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (amd64, gcc12, openmpi) / Dev environment (Python)	✅ success	view
Documentation / Build and check docs	✅ success	view

[1tnguyen]

LGTM 👍 only a couple of nits

[Renaud-K]

/ok to test edfb8ad

Command Bot: Processing...

[Renaud-K]

I have also run an independent review on this PR and applied the following feedback.

• P2 — kernel lookup misclassifies every driver failure as an unknown name.
  get_kernel failures all return CUDAQ_ERR_INVALID_ARG. Errors such as invalid context/handle or driver deinitialization are reported as “no kernel.” Match handler lookup behavior: treat only CUDA_ERROR_NOT_FOUND as invalid input; report other CUresults and return CUDAQ_ERR_CUDA.

• P1
  If a launcher enqueues the persistent kernel and then returns an error, start() destroys and clears the stream without stopping or synchronizing it. Later destruction sees a null stream and may unload the linked module while its kernel is active.
  Either explicitly require that non-OK means “nothing was enqueued,” or track possible launch activity and synchronize before unloading. Add a launcher test that enqueues successfully, then returns failure.

[github-actions]

CI Summary (merge_group) — ✅ passed

Run #40996302579 · ✅ 5 · ⏩ 13 · ❌ 0 · ⛔ 0

⚡ Fast merge queue — only the llvm dev-environment build+test legs ran.
Wheels, metapackages, installers, Docker images, and the gcc12 legs were skipped.
Set the repository variable MERGE_QUEUE_FAST to false to restore the full matrix.

Top-level jobs (18)

Job	Result
binaries	⏩ skipped
build_and_test	✅ success
changes	✅ success
config_devdeps	✅ success
config_source_build	⏩ skipped
config_wheeldeps	⏩ skipped
devdeps	✅ success
docker_image	⏩ skipped
documentation	⏩ skipped
gen_code_coverage	⏩ skipped
metadata	✅ success
python_devel_wheel	⏩ skipped
python_metapackages	⏩ skipped
python_wheels	⏩ skipped
python_wheels_cu12_validation	⏩ skipped
python_wheels_cu13	⏩ skipped
source_build	⏩ skipped
wheeldeps	⏩ skipped

⏩ Skipped jobs (13) — intentionally skipped on PR builds; run on merge_group / workflow_dispatch

Job
binaries
config_source_build
config_wheeldeps
docker_image
documentation
gen_code_coverage
python_devel_wheel
python_metapackages
python_wheels_cu12_validation
python_wheels_cu13
python_wheels
source_build
wheeldeps

All sub-jobs (28) — every matrix leg, with links

Job	Status	Link
Build and test (amd64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (amd64, llvm, openmpi) / Dev environment (Python)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Python)	✅ success	view
CI Summary	❔ in_progress	view
Check for stable CUDA-Q changes	✅ success	view
Configure build (devdeps)	✅ success	view
Configure build (source_build)	⏩ skipped	view
Configure build (wheeldeps)	⏩ skipped	view
Create CUDA 12 Python wheels	⏩ skipped	view
Create CUDA 13 Python wheels	⏩ skipped	view
Create CUDA Quantum installer	⏩ skipped	view
Create Docker images	⏩ skipped	view
Create Python metapackages	⏩ skipped	view
Create shared Python development wheel	⏩ skipped	view
Documentation	⏩ skipped	view
Gen code coverage	⏩ skipped	view
Load dependencies (amd64, llvm) / Caching	✅ success	view
Load dependencies (amd64, llvm) / Finalize	✅ success	view
Load dependencies (amd64, llvm) / Metadata	✅ success	view
Load dependencies (arm64, llvm) / Caching	✅ success	view
Load dependencies (arm64, llvm) / Finalize	✅ success	view
Load dependencies (arm64, llvm) / Metadata	✅ success	view
Load source build cache	⏩ skipped	view
Load wheel dependencies	⏩ skipped	view
Prepare cache clean-up	❔ in_progress	view
Retrieve PR info	✅ success	view
Validate CUDA 12 Python wheels	⏩ skipped	view

✅ Required checks (4/4) — declared in .github/required-checks.yml for merge_group_fast

Required check	Status	Link
Build and test (amd64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (amd64, llvm, openmpi) / Dev environment (Python)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Debug)	✅ success	view
Build and test (arm64, llvm, openmpi) / Dev environment (Python)	✅ success	view