CUDA: pass descriptor-table entry-point uniforms by reference (split from #11939; supersedes #11941)

[szihs]

This is the by-reference layer of the #11774 fix, split back out of #11939 and stacked on it. Base branch = haaggarwal/cuda-param-dynamic-index-floor (#11939, the ABI-invisible floor); this PR is only the breaking ABI change. It supersedes #11941, which was previously merged into #11939's branch and cannot be reopened (GitHub blocks reopening merged PRs). Review this against #11939 as its base — the diff shown is by-ref only.

Fixes #11774 (with #11939).

Motivation

#11939's local-copy legalization is the floor: it makes runtime-indexed by-value parameters correct and O(1)-per-access, but the host still ships the full descriptor blob by value per dispatch and every thread pays a prologue copy. The measured regression — SlangPy test_tensor_sum_indirect, 13–29× slower than Vulkan on L40S — is a uniform TensorList<N> struct nesting RWTensor<float,2>[N] in .param. On every other target, indexed resource collections live in dynamically-indexable descriptor memory; this PR gives CUDA the same placement.

Proposed solution

A compute entry-point uniform struct parameter that carries a descriptor table (a fixed-size array whose elements are or contain resources/pointers) is passed by reference — laid out and reflected as an implicit ParameterBlock, so the kernel receives one 8-byte device pointer and the payload lives in pooled device global memory:

__global__ void main(TensorList_0* list_0, ...)   // was: TensorList_0 list_0 (by value)

The decision is made once, at parameter binding (shouldPassCUDAEntryPointUniformParamByRef), where reflection is computed; the post-link IR pass reconcileCUDAByRefEntryPointParams consumes the recorded layout rather than re-deriving a predicate — so the emitted signature and the reflected layout cannot disagree, eliminating by construction the ABI-desync class that sank the earlier __constant__-hoist attempt (#11747). Zero emit changes (ParameterBlock params already render as T*), zero slang-rhi changes (entry-point-parented PB sub-objects are an existing, tested path). -cuda-entry-point-params-by-value restores the legacy ABI.

Scope: struct-typed params only (bare descriptor arrays would reflect as ParameterBlock<array>, a shape hosts don't yet represent — they keep the by-value layout and rely on #11939's floor). Excluded: RT stages, torch/[AutoPyBindCUDA]/[CudaKernel] kernels, explicit parameter groups, plain-data params.

Change summary

Area	Files
Decision site	slang-parameter-binding.cpp
IR reconciliation	slang-ir-cuda-byref-entry-point-params.{h,cpp}; rewriteValueUsesToAddrUses hoisted to slang-ir-transform-params-to-constref.{h,cpp}; invocation in slang-emit.cpp
Option	include/slang.h (CudaEntryPointParamsByValue), slang-options.cpp, docs/command-line-slangc-reference.md
Docs	docs/cuda-target.md (Binding rewritten), docs/user-guide/09-targets.md
Tests	tests/cuda/entry-point-uniform-*.slang — emit + reflection-JSON + opt-out + exclusions + a GPU execution test

Validation

All perf/nsys numbers on record were measured with this PR + #11939 applied together (floor + by-ref), since the two land as a stack. Runtime-indexed descriptor tables 13–29× faster on L40S + RTX 5090, CUDA beating Vulkan (0.173 vs 0.199 ms at count=32); command stream adds exactly one async pooled cuMemcpyHtoDAsync per command buffer (nsys A/B 31:31:0 vs by-value control); slangpy CUDA canaries green with shader-slang/slangpy#1045. All 53 tests/cuda pass on the stacked branch (46 on #11939 alone).

Landing order

Label pr: breaking change. Merge order: #11939 (non-breaking floor) → shader-slang/slangpy#1045 (host reference-safety, before slangpy bumps its Slang pin) → this PR → close #11774.

🤖 Generated with Claude Code

[coderabbitai]

Important

Review skipped

Draft detected.

Please check the settings in the CodeRabbit UI or the .coderabbit.yaml file in this repository. To trigger a single review, invoke the @coderabbitai review command.

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: ASSERTIVE

Plan: Pro

Run ID: df004e6f-8c87-435c-a2e3-02fd0830b8d4

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

• 🔍 Trigger review

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}

[szihs]

This supersedes #11941 (which couldn't be reopened)

#11941 was previously merged into #11939's branch, and GitHub does not allow reopening a merged PR — so the by-reference layer is re-proposed here as a fresh PR, stacked on #11939 (base = haaggarwal/cuda-param-dynamic-index-floor). The diff shown is by-ref only; review it against #11939 as the base.

Provenance of the results: all perf, nsys command-stream, and correctness numbers on record (#11774 and #11939) were measured with #11939 (floor) + this PR (by-ref) applied together — they were developed and validated as one stack and land as one. Specifically attributable to this PR (the memory-placement change):

• runtime-indexed descriptor tables 13–29× faster (L40S + RTX 5090), CUDA beating Vulkan at 0.173 vs 0.199 ms (count=32);
• exactly one async pooled cuMemcpyHtoDAsync per command buffer (nsys A/B 31:31:0 vs a by-value control) — verified async into fresh per-encode memory, deliberately not the rejected per-dispatch __constant__ symbol copy of [Draft] Hoist CUDA runtime-indexed resource-array entry-point params to a parameter group #11747;
• 53/53 tests/cuda on the stacked branch (vs 46/46 on CUDA: legalize dynamic indexing into by-value kernel parameters (floor for #11774) #11939 alone — the delta is this PR's reflection/emit/exclusion/execution tests).

Landing dependency: shader-slang/slangpy#1045 (host reference-safety) must merge before slangpy bumps its Slang pin — this PR changes the CUDA ABI that #1045 makes slangpy's cached-offset writers safe for. That is why the combined PR's SlangPy Tests check was red; it is gated on #1045, not on a defect here.

The split is faithful: #11939 + this PR reproduce byte-for-byte the combined tree that was tested.

[szihs]

Command-stream evidence for this by-reference layer (originally posted on the pre-split #11939; measured with floor + by-ref together).

nsys A/B: command-stream composition with a by-ref-converted parameter (L40S, measured)

Question: a shader that previously dispatched with a single CUDA command-stream entry gets one entry-point parameter converted to by-reference — is an extra cuMemcpy now needed?

Answer: yes — exactly one cuMemcpyHtoDAsync per command buffer (not per dispatch beyond that, and no sync copy). Measured on L40S with the same slangpy kernel (Array1DValueType<Val,4> — the converting descriptor-table shape), 31 dispatches, A/B via a one-line gate toggle with md5-verified library swaps:

CUDA API call	A (by-ref)	B (by-value control)
cuLaunchKernel	31	31
cuMemcpyHtoDAsync	31	0
cuMemcpy (inputs + readback, unchanged)	66	66

GPU-trace memops corroborate: [CUDA memcpy HtoD] delta is exactly 31 — one payload upload per command buffer. slangpy submits one command buffer per call, hence 31:31; a program batching N dispatches per command buffer sees 1 upload for N launches (the upload is issued once in CommandExecutor::execute, cuda-command.cpp:293, before the command-replay loop).

Properties of that copy, matching the source line-by-line:

• Async into fresh per-encode pool memory (ConstantBufferPool::upload, cuda-constant-buffer-pool.cpp:23-54) — no sync point, pipelines with the launch, and no write-after-read hazard between in-flight command buffers. Deliberately not the rejected per-dispatch cuMemcpyToSymbol-into-__constant__ design ([Draft] Hoist CUDA runtime-indexed resource-array entry-point params to a parameter group #11747), which serialized back-to-back dispatches on a module symbol.
• The payload goes to the global pool ("Sub-objects are always written to global memory, even if the parent represents an entry-point", cuda-shader-object.cpp); the entry-point blob stays host-only and reaches the kernel via CU_LAUNCH_PARAM_BUFFER_POINTER as before, now 8 bytes smaller per converted param than the inline payload it replaces.
• If the program already uses any global-scope parameter or parameter block, that pooled upload already exists and the converted payload rides it — zero additional stream entries.
• The end-to-end perf numbers in the PR description already include this copy; it only applies to descriptor-table-carrying params (10-30x degenerate by-value), and -cuda-entry-point-params-by-value restores the single-entry stream wholesale.

Repro for an independent Nsight session: slangpy/benchmarks/profile_byref_pool_upload.py on the slangpy branch haaggarwal/param-array-copy-benchmark (--mode separate vs --mode batched shows the per-command-buffer, not per-dispatch, nature directly).

[szihs]

Scaling follow-up to the above (originally on #11939 pre-split).

Follow-up to the nsys A/B above — scaling measurement (L40S) confirming the upload is per command buffer, not per dispatch:

mode / dispatches	cuLaunchKernel	cuMemcpyHtoDAsync
separate submits / 100	101	101
batched (one command buffer) / 100	101	3
batched (one command buffer) / 200	201	5

Separate submits scale 1:1 (the honest per-dispatch cost in slangpy's one-command-buffer-per-call pattern). Batched stays flat while launches double — the residual 3→5 is one copy per used global-pool page (verified via the GPU memop trace, transfers up to ~12 KB), i.e. the count tracks payload volume, not dispatch count. Repro: slangpy/benchmarks/profile_byref_pool_upload.py (slangpy branch haaggarwal/param-array-copy-benchmark), which also now includes a launch-bound benchmark of the worst case — a statically indexed descriptor table, where the upload is pure per-call overhead — for putting a hard number on that fixed cost A/B against the legacy ABI.