Skip to content

Introduce the target agnostic UNIFIED kernel architecture with nvJitLink - #5722

Merged
Renaud-K merged 9 commits into
NVIDIA:mainfrom
Renaud-K:gpu-dataplane
Oct 29, 2026
Merged

Introduce the target agnostic UNIFIED kernel architecture with nvJitLink#5722
Renaud-K merged 9 commits into
NVIDIA:mainfrom
Renaud-K:gpu-dataplane

Conversation

@Renaud-K

@Renaud-K Renaud-K commented Oct 24, 2026 •
edited
Loading

Copy link
Copy Markdown
Collaborator

Summary

Until now the DEVICE + UNIFIED dispatch kernel was built for one transport at
compile time: libcudaq-realtime-dispatch.a compiled the kernel together with
the transport's device hooks and every application device-linked that archive with its
handlers. A binary could only ever drive the one transport it was built
against.

This PR makes the unified kernel target-agnostic. The kernel now ships once,
as LTO-IR embedded in libcudaq-realtime.so. At run time, the dispatcher links
three pieces into one CUDA module with nvJitLink, for the GPU it runs on:

  • the kernel,
  • the device hooks of whichever bridge cudaq_dispatcher_set_bridge wired,
  • the application's DEVICE_CALL handlers.

A device data plane is now selected at run time, like any other bridge.

flowchart LR
  kernel["Unified kernel: LTO-IR embedded in libcudaq-realtime.so"]
  bridge["Bridge hooks: -rdc device code in the bridge library"]
  app["DEVICE_CALL handlers: -rdc device code in the application"]
  link["nvJitLink -arch=sm_XY -lto, on the first handler lookup"]
  module["One CUDA module, run by the dispatcher"]
  kernel --> link
  bridge -->|"get_gpu_dataplane"| link
  app -->|"cudaq_dispatcher_get_device_handler"| link
  link --> module
Loading

Changes by part

Application-facing API

File Role
include/cudaq/realtime/daemon/dispatcher/cudaq_realtime.h Adds cudaq_dispatcher_get_device_handler, the one call an application makes for the device link: it returns the address of an exported handler in the dispatcher's linked module, and the first call performs the link. A header-only static C function, so it runs in, and reads the device code of, the binary that holds the handler. It forwards to the exported cudaq_dispatcher_get_device_handler_in. Removes cudaq_launch_unified_dispatch_device. Documents the DEVICE + UNIFIED precedence of set_bridge.
include/cudaq/realtime/daemon/dispatcher/cudaq_status.h (new) cudaq_status_t, moved out of cudaq_realtime.h as a first step so that bridges do not need to depend on the dispatcher API. (More work is needed)
include/cudaq/realtime/device_code/device_code.h (new) cudaq_device_image_t (relocatable device code: one fatbin, or several back to back as in a binary's __nv_relfatbin section) and CUDAQ_REALTIME_EXPORT_DEVICE_FUNCTION, which exports a device function through a pointer global, since CUDA looks up kernels and globals by name but never device functions.
include/cudaq/realtime/device_code/device_code_reader.h (new) cudaq_device_code_of_this_binary: a few lines of header-only C that return the calling binary's __nv_relfatbin section as one image. It has to be compiled into each binary because the hidden __start___nv_relfatbin and __stop___nv_relfatbin symbols resolve to the binary that references them.

Bridge interface

File Role
include/cudaq/realtime/daemon/bridge/bridge_interface.h Interface version 4: cudaq_gpu_dataplane_t { ctx, device_code, device_code_count } and a get_gpu_dataplane vtable entry, plus the public cudaq_bridge_get_gpu_dataplane, mirroring cudaq_bridge_get_cpu_dataplane.
include/cudaq/realtime/daemon/dispatcher/unified_device_transport.cuh The device hook contract: cudaq_dev_enter_kernel, cudaq_dev_rx_poll, cudaq_dev_tx_publish, and the session storage size. The LINKAGE note now describes the run-time link.
lib/daemon/bridge/bridge_interface_api.cpp cudaq_bridge_get_gpu_dataplane: CUDAQ_ERR_UNSUPPORTED for a provider that leaves the entry NULL.
lib/daemon/bridge/udp/bridge_impl.cpp, lib/daemon/bridge/cpu_roce/bridge_impl.cpp A NULL get_gpu_dataplane entry; these providers have no GPU data plane.

Dispatcher

File Role
lib/daemon/dispatcher/cudaq_realtime_api.cpp set_bridge, DEVICE + UNIFIED: a launch override (the UNIFIED transport context's launch_fn) takes precedence; otherwise the bridge's GPU data plane, whose images are copied; a bridge with neither is refused. Link: link_dispatcher links the embedded kernel, the bridge's images and the application's images, and finds the kernel by name. Lookup: cudaq_dispatcher_get_device_handler_in links on the first call, reuses the module afterwards, refuses device code from another binary (compared by content, so every translation unit of one binary works), and records each address it hands out. start: links the kernel and the hooks alone if no handler was looked up, refuses any DEVICE_CALL entry whose address did not come from the lookup, and launches the linked kernel with cudaLaunchKernel. destroy: leaves the module loaded while a kernel may still run from it. The JIT state is grouped in JitDeviceLibrary.
lib/daemon/dispatcher/unified_dispatch_kernel.cu The kernel itself, now extern "C" so it can be found by name in the linked module, and compiled to LTO-IR only. It has no transport dependency.
lib/daemon/dispatcher/unified_transport_none.cu (deleted), lib/daemon/CMakeLists.txt The compile-time transport selection, the null transport and the unified kernel's build in libcudaq-realtime-dispatch.a are gone.

Run-time device link (lib/device_code/, new)

A static library, libcudaq-realtime-device-code.a, linked privately into libcudaq-realtime.so.

File Role
CMakeLists.txt Builds the kernel's LTO-IR fatbin with a custom nvcc command, at CUDAQ_REALTIME_DEVICE_CODE_MIN_ARCH (default 80). LTO-IR for one architecture links for that GPU and every newer one, and the kernel has no __CUDA_ARCH__ branches. Links CUDA::nvJitLink PRIVATE.
unified_dispatch_kernel_image.cpp Embeds the fatbin between two hidden symbols with asm(".incbin ..."). It uses only generic GNU assembler directives, so it is portable across ELF targets (x86-64, AArch64).
nvjitlink_linker.{h,cpp} link_device_code: splits each image into its fatbins, adds them to nvJitLink, links with -arch=sm_<this GPU> -lto, and loads the cubin with cudaLibraryLoadData. Failures print the nvJitLink version and its error log.
lib/CMakeLists.txt Adds the directory.

gpu_roce bridge

File Role
lib/daemon/bridge/gpu_roce/bridge_impl.cpp Implements get_gpu_dataplane in unified mode: stages the DOCA transport context in device memory once per bridge, and hands over the library's device code. Refuses UNIFIED (no launch override) and RING_BUFFER in unified mode.
lib/daemon/bridge/gpu_roce/gpu_roce_unified_transport.cu, gpu_roce_device_transport.cuh (new) The GpuRoceTransceiver/DOCA device transport: CQE polling, receive recycling, BlueFlame or proxied send. It includes no CUDA-Q realtime header and uses its own names (gpu_roce_dev_*), status enum and session size, so it can move to hololink as is.
lib/daemon/bridge/gpu_roce/gpu_roce_device_hooks.cu (new) The bridge's cudaq_dev_* hooks: one-line forwarders to the transport, which map its status and check that its session fits the kernel's storage. Only the bridge sees the fixed hook names.
lib/daemon/bridge/gpu_roce/CMakeLists.txt Builds both device files with -rdc and one LTO-IR entry per architecture in CUDA_TARGET_ARCHS: DOCA picks its fences by __CUDA_ARCH__. With LTO-IR, the run-time link inlines transport, forwarder and kernel.
include/cudaq/realtime/gpu_roce_bridge_common.h bridge_run no longer drives the unified kernel; --unified refers to gpu_roce_app.

Tests and tools

File Role
unittests/bridge_interface/gpu_loopback/ (new) gpu_loopback_bridge.cu is a test-only provider with a GPU data plane over mapped pinned memory, so the whole run-time link runs on any GPU without a NIC. test_gpu_dataplane.cpp covers: serving requests; the address and binary guards; relinking after a failed link; a provider image that cannot link; a provider without device code; the reader in a binary without -rdc code (no_device_code.cpp); and a binary with two fatbins whose second handler is looked up from its own translation unit (test_gpu_dataplane_second_handler.cu).
unittests/c_api/ (new) c_api_check.c compiles the public headers and the header-only wrapper as C. check_no_nvjitlink.cmake fails if a public header includes nvJitLink.h.
unittests/test_server/cudaq_realtime_test_server.cpp, unittests/bridge_interface/gpu_roce/gpu_roce_bridge.cpp Unified wiring through cudaq_dispatcher_get_device_handler.
unittests/utils/rpc_increment_handler.{cu,cuh} (new), init_rpc_increment_function_table.cu, CMakeLists.txt The increment handler, defined once and exported, compiled by each executable as an INTERFACE source. The linked-table setup looks the handler up itself.
unittests/utils/gpu_roce_test.sh --unified runs gpu_roce_app and waits for its dispatcher to start.
unittests/CMakeLists.txt Adds the new test directories.
scripts/install_devdeps.sh Installs libnvjitlink-dev, needed to build libcudaq-realtime only.

Documentation

File Role
realtime/docs/device_code_contract.md (new) The contract the bridge and the application follow for the run-time link (summarized below).
realtime/docs/cudaq_realtime_host_api.md, docs/sphinx/using/realtime/host.md The unified wiring sequence, cudaq_dispatcher_get_device_handler and the set_bridge shape table.
realtime/docs/cudaq_realtime_network_interface.md get_gpu_dataplane: how a provider serves the unified kernel.

Requirements for application developers

  • Build the handlers as relocatable device code: -rdc=true, or CUDA_SEPARABLE_COMPILATION ON.
  • Export each handler with CUDAQ_REALTIME_EXPORT_DEVICE_FUNCTION(name) (device_code.h). The handler itself may have internal linkage.
  • Compile the export into the binary that runs the lookup, not into a static-library member: nothing on the host references it, so the linker would drop it. A library can carry it as an INTERFACE source.
  • Look the handler up after cudaq_dispatcher_set_bridge with cudaq_dispatcher_get_device_handler(dispatcher, "name", &fn), and put the returned address in the DEVICE_CALL entry. An address taken in the application's own module (for example by an init kernel) is refused by start.
  • One binary per dispatcher: all of a dispatcher's handlers come from one binary (an executable or a shared library), and are looked up before start.
  • No host-set device state in handlers yet: a __device__ global set from the host lives in the application's module, not in the linked one.
  • Run time: libnvJitLink.so.<major>, which ships with the CUDA runtime, at least as new as every toolkit that produced the linked code. Applications do not need the nvJitLink dev package.
  • No longer needed for unified dispatch: linking libcudaq-realtime-dispatch.a, or calling cudaq_launch_unified_dispatch_device.
  • Optional: handler code that branches on __CUDA_ARCH__ should ship an entry per architecture it cares about.

Requirements for bridge developers

  • Implement the three hooks of unified_device_transport.cuh. The in-tree pattern (gpu_roce) keeps the transport's device code free of CUDA-Q headers, with its own names, and defines cudaq_dev_* in the bridge as one-line forwarders. The session must fit CUDAQ_DEV_TRANSPORT_SESSION_BYTES and CUDAQ_DEV_TRANSPORT_SESSION_ALIGN.
  • Build the hooks with -rdc. Add LTO-IR so the run-time link inlines them into the kernel (--generate-code=arch=compute_XX,code=lto_XX per source, to avoid host LTO), with one entry per architecture if the code branches on __CUDA_ARCH__, as DOCA's headers do. Without LTO-IR the link still works, but the hooks stay real calls on every request.
  • Implement get_gpu_dataplane (interface version 4). Return a device-accessible ctx and the image from cudaq_device_code_of_this_binary. The image storage must outlive the bridge, so keep it in the bridge's context.
  • Keep every other device symbol private (static or an anonymous namespace), including those defined in headers: the link takes all of the binary's -rdc code, and two definitions of one name fail it.
  • Link no realtime library: a provider uses only the header-only code. libcudaq-realtime.so loads the provider, never the reverse.
  • Rebuild against version 4: positional vtable initializers need a get_gpu_dataplane entry, NULL for a provider with no GPU data plane.
  • Launch override: only a transport that replaces the whole dispatch loop answers the UNIFIED transport context with a launch_fn. Others refuse it, or answer with a NULL launch_fn, and set_bridge then uses the GPU data plane.
  • Toolkit: images must come from a CUDA toolkit no newer than the nvJitLink the application runs with.

Breaking changes

  • cudaq_launch_unified_dispatch_device and the unified kernel in libcudaq-realtime-dispatch.a are removed, together with the CUDAQ_REALTIME_UNIFIED_TRANSPORT_TU property and the null transport.
  • DEVICE_CALL entries of a DEVICE + UNIFIED dispatcher must come from cudaq_dispatcher_get_device_handler.
  • The bridge interface is now version 4, so providers must rebuild.
  • cudaq_dispatcher_set_unified_launch now requires a launch function.
  • bridge_run --unified is removed; use gpu_roce_app --unified.

Testing

  • Realtime unit tests: 109/109 ctest on GB200, including the new GpuDataplaneTest suite and the C compile check. The GPU data plane tests skip without a GPU.
  • RoCE two-process tests: 30/30 ctest -R TwoProcess on RoCE hardware (ConnectX, DOCA 3.4, Hololink 2.7), ring and unified shapes, for udp, cpu_roce and gpu_roce.
  • End to end against the FPGA: gpu_roce_test.sh --unified on GB200 GPU 3 (on the NIC's NUMA node), 100 requests per run, 5 runs, all 100/100 matched. The average round trip is 5.94 µs (median of runs), the same as with the hooks compiled into the kernel (5.96 µs).
  • Inlining matters: an experiment calling the hooks through function pointers measured 6.56 µs, about 10% slower, so the bridge forwarders are inlined, not called.

@copy-pr-bot

copy-pr-bot Bot commented Oct 24, 2026

Copy link
Copy Markdown

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.

@github-actions github-actions Bot added documentation Improvements or additions to documentation build The issue related to the CUDA Quantum build realtime Related to CUDA-Q realtime labels Oct 24, 2026
@Renaud-K
Renaud-K requested review from 1tnguyen, bettinaheim, boschmitt and cketcham2333 and removed request for cketcham2333 October 24, 2026 01:47
@Renaud-K

Renaud-K commented Oct 24, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test 40215d0

Command Bot: Processing...

@github-actions

github-actions Bot commented Oct 24, 2026 •
edited
Loading

Copy link
Copy Markdown

CI Summary (push) — ✅ passed

Run #40808310697 · ✅ 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

@Renaud-K

Renaud-K commented Oct 25, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test dcd9ed1

Command Bot: Processing...

@Renaud-K

Renaud-K commented Oct 26, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test ff5d2c7

Command Bot: Processing...

@Renaud-K

Renaud-K commented Oct 26, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test 98011c2

Command Bot: Processing...

@Renaud-K

Renaud-K commented Oct 26, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test 0fa37a7

Command Bot: Processing...

@Renaud-K
Renaud-K requested a review from mitchdz October 27, 2026 16:34
Renaud-K
Renaud-K commented Oct 27, 2026
View reviewed changes
Comment thread realtime/include/cudaq/realtime/daemon/bridge/bridge_interface.h
1tnguyen
1tnguyen reviewed Oct 28, 2026
View reviewed changes
Comment thread realtime/include/cudaq/realtime/device_code/device_code_reader.h Outdated
1tnguyen
1tnguyen reviewed Oct 28, 2026
View reviewed changes
Comment thread realtime/include/cudaq/realtime/daemon/dispatcher/cudaq_realtime.h Outdated
@1tnguyen

1tnguyen commented Oct 28, 2026

Copy link
Copy Markdown
Collaborator

LGTM 👍

@Renaud-K
gpu_dataplane
7c90add 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K
realtime: keep nvJitLink private to libcudaq-realtime
8dd5fb1 
Link CUDA::nvJitLink PRIVATE into cudaq-realtime-device-code, add a ctest
that fails if a public header includes nvJitLink.h, and document that only
building libcudaq-realtime needs the nvJitLink dev package while applications,
bridges and customers need only the runtime libnvJitLink.so.<major>.
@Renaud-K
Introducing inlined wrappers to eliminate CUDAQ dependency
0d795ac 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K
Fixing ninja version
1ede00d 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K
Adding support for CUDA 12.6
44d432e 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K
Fixing CI
e6fffe9 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K
Merging in PR 5675
ade2e24 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K
Review comments
121d3a8 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K

Renaud-K commented Oct 28, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test 121d3a8

Command Bot: Processing...

@Renaud-K
Applied review feedback
82a534d 
Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
@Renaud-K

Renaud-K commented Oct 29, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Collaborator Author

/ok to test 82a534d

Command Bot: Processing...

@Renaud-K
Renaud-K added this pull request to the merge queue Oct 29, 2026
@github-actions

github-actions Bot commented Oct 29, 2026

Copy link
Copy Markdown

CI Summary (merge_group) — ✅ passed

Run #40927289187 · ✅ 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 ✅ success 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

Merged via the queue into NVIDIA:main with commit e78671b Oct 29, 2026
108 of 109 checks passed
@Renaud-K
Renaud-K deleted the gpu-dataplane branch October 29, 2026 19:33
github-actions Bot pushed a commit that referenced this pull request Oct 29, 2026
@cuda-quantum-bot
Cleaning up docs preview for PR #5722.
a8bffcb
@Renaud-K Renaud-K mentioned this pull request Oct 29, 2026
Merged
sacpis pushed a commit to sacpis/cuda-quantum that referenced this pull request Oct 30, 2026
@Renaud-K
Adding support for a custom GPU dispatcher (NVIDIA#5778)
a46d879 
# 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 NVIDIA#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:

```cmake
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`:

```cpp
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

```mermaid
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
```

**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 NVIDIA#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.

---------

Signed-off-by: Renaud Kauffmann <rkauffmann@nvidia.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@1tnguyen 1tnguyen 1tnguyen left review comments
@bettinaheim bettinaheim bettinaheim approved these changes
@cketcham2333 cketcham2333 Awaiting requested review from cketcham2333
@boschmitt boschmitt Awaiting requested review from boschmitt
@mitchdz mitchdz Awaiting requested review from mitchdz

Assignees

No one assigned

Labels

build The issue related to the CUDA Quantum build documentation Improvements or additions to documentation realtime Related to CUDA-Q realtime

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@Renaud-K @1tnguyen @bettinaheim