Commit 0d8e757
## Motivation
There is no public way to ask Slang which version of a downstream
pass-through compiler it will
*actually load and use*. This matters for NVRTC: an application (Falcor)
wants to guard a CUDA
diagnostic path that is pathological on older NVRTC and fine on newer,
and the guard must key off
the NVRTC library *Slang* selected — Slang's discovery order (an
explicit
`setDownstreamCompilerPath`, then the instance directory, then
`CUDA_PATH`, then `PATH`, picking the
newest match) can differ from a version the application discovers on its
own by calling
`nvrtcVersion()`.
Concrete consumer:
```cpp
int major = 0, minor = 0;
if (SLANG_SUCCEEDED(globalSession->getDownstreamCompilerVersion(SLANG_PASS_THROUGH_NVRTC, &major, &minor)))
useOldNvrtcWorkaround = (major < 12) || (major == 12 && minor < 7);
```
The version is already captured internally:
`NVRTCDownstreamCompiler::init()` calls the loaded
library's `nvrtcVersion(&major, &minor)` and stores it in
`m_desc.version`
(`source/compiler-core/slang-nvrtc-compiler.cpp:196-197`). This change
just exposes that captured
value through a public, ABI-additive method.
## Proposed solution
Append a single method to `IGlobalSession`:
```cpp
virtual SLANG_NO_THROW SlangResult SLANG_MCALL
getDownstreamCompilerVersion(SlangPassThrough passThrough, int* outMajor, int* outMinor) = 0;
```
The implementation routes through
`Session::getOrLoadDownstreamCompiler(passThrough, nullptr)` — the
same memoized lazy-discovery funnel that compilation uses — and reads
the numeric version from
`getDesc().version`. Routing through that funnel is the load-bearing
correctness property: it shares
the memoized compiler cache with `SLANG_PTX` compilation and honors
`setDownstreamCompilerPath` plus
the standard search order, so the reported version is guaranteed to be
the library Slang will
actually compile with.
**Why numeric (`int* outMajor, int* outMinor`) and not a version
string?** A string-blob shape
(`getVersionString`) is unusable for the very compiler this feature
targets: `NVRTCDownstreamCompiler`
does **not** override `getVersionString` (only DXC, glslang, and Tint
do), so it inherits the
`DownstreamCompilerBase` default that returns `SLANG_FAIL` and a null
blob
(`source/compiler-core/slang-downstream-compiler.h:401-406`). The
version is reliably reachable only
through `getDesc().version`. A numeric pair is also directly usable for
the version-comparison guard
that is the actual use case — no string-format contract to freeze and no
client-side parsing. A
public descriptor struct was rejected as an unnecessary permanent ABI
commitment.
> **Maintainer decision point (public ABI shape).** This is a permanent
public-API surface. We
> recommend the numeric `int* outMajor, int* outMinor` shape above for
the reasons given. If you
> prefer a different shape — e.g. a version *string* blob (would
additionally require adding a
> `getVersionString` override to `NVRTCDownstreamCompiler`), or a richer
descriptor — we are happy to
> adjust. Flagging because the choice is hard to change once released.
## Change summary
| File | Change |
|------|--------|
| `include/slang.h` | Append the new pure-virtual to `IGlobalSession`,
immediately before the closing `};` (after `saveBuiltinModule`) —
append-only, ABI-safe. |
| `source/slang/slang-global-session.h` | Declare the `Session`
override. |
| `source/slang/slang-global-session.cpp` | Implement: boundary-validate
the pass-through, load via `getOrLoadDownstreamCompiler`, read
`getDesc().version`. |
| `source/slang-record-replay/proxy/proxy-global-session.h` |
`GlobalSessionProxy` override (record/replay forwarding, records the
version out-params). |
| `source/slang-record-replay/replay-handlers.cpp` | Register the new
method's replay handler so recorded streams containing the call dispatch
on replay. |
| `tools/slang-unit-test/unit-test-vtable-stability.cpp` | Add the new
method to the `IGlobalSession` vtable-layout probe (slot 32) and assert
it. |
| `tools/slang-unit-test/unit-test-downstream-compiler-version.cpp`
(new) | GPU-free unit test for the new API. |
## Concepts and vocabulary
- **`getOrLoadDownstreamCompiler`** — the single, mutex-guarded,
memoized funnel that finds and loads
a downstream compiler for a pass-through, applying the
override/instance-dir/`CUDA_PATH`/`PATH`
search order and caching the result. All compilation goes through it;
routing the query through it
is what guarantees the reported version matches the library used for
codegen.
- **`DownstreamCompilerDesc::version`** — a `SemanticVersion`
(`m_major`/`m_minor`/`m_patch`) on the
desc returned by `IDownstreamCompiler::getDesc()`. For NVRTC it is
populated from `nvrtcVersion`
at load time; for compilers that report no version it is `(0,0)`.
- **vtable-stability probe** —
`tools/slang-unit-test/unit-test-vtable-stability.cpp` pins the COM
vtable slot index of every `IGlobalSession` method; appending a method
adds the next slot, which
the test now covers, guarding against accidental mid-vtable insertion.
- **record/replay proxy** — `GlobalSessionProxy` wraps a real
`IGlobalSession` to record API calls
for later replay; every interface method needs a forwarding override
here.
## Process report
- **`include/slang.h` (append the virtual).** The method is appended at
the very end of
`IGlobalSession`, after the previously-last method `saveBuiltinModule`,
with an explicit doc
comment. This is the only ABI-safe placement: inserting mid-interface
would shift every subsequent
vtable slot and break callers compiled against the old header. No enum
or struct is added, so the
change is `pr: non-breaking`.
- **`source/slang/slang-global-session.h` (declaration).** Declares the
override on the concrete
`Session` class. This is required because `Session` implements
`IGlobalSession`; an appended pure
virtual makes the class abstract until overridden. The declaration order
here is cosmetic — the
vtable slot order is fixed solely by the interface declaration in
`include/slang.h`.
- **`source/slang/slang-global-session.cpp` (implementation).** Mirrors
the existing
`checkPassThroughSupport` → `getOrLoadDownstreamCompiler` precedent for
load and error semantics
(`SLANG_E_NOT_FOUND` when the compiler is not loadable). It then reads
the numeric version directly
from `compiler->getDesc().version` rather than from `getVersionString`.
*Input-shape check:* the data read here — `getDesc().version` — is the
correct, principled source.
It is exactly the field `NVRTCDownstreamCompiler::init()` populates from
the loaded library and the
same field NVRTC uses to gate its own behavior, so there is no upstream
producer to fix; the
string accessor is simply not implemented for NVRTC, which is why we
read the desc. A public
boundary check rejects out-of-range `SlangPassThrough` values up front
(the loader indexes
per-type arrays by the enum value, so an out-of-range value would
otherwise index out of bounds);
`SLANG_PASS_THROUGH_NONE` returns `SLANG_E_NOT_FOUND` because there is
no compiler to report (this
intentionally differs from `checkPassThroughSupport(NONE) == SLANG_OK`).
`outMajor`/`outMinor` are
optional (null tolerated). A loadable-but-versionless compiler yields
`SLANG_OK` with `(0,0)`,
consistent with "the loaded compiler's reported version."
- **`source/slang-record-replay/proxy/proxy-global-session.h` (proxy
override).** Required because
`GlobalSessionProxy` implements `IGlobalSession`; without the override
the class would not compile.
It records the call and input, then — following the scalar-out-param
pattern already used by
`getTypeConformanceWitnessSequentialID` in the session proxy — uses
`PREPARE_POINTER_OUTPUT` (which
redirects a null arg to a stack temp, making the subsequent record
null-safe) and `RECORD_OUTPUT`
so the deterministic version values are captured for faithful replay.
- **`source/slang-record-replay/replay-handlers.cpp` (replay
registration).** Every `IGlobalSession`
method that the proxy records has a matching `REPLAY_REGISTER` entry;
without one, a recorded
stream containing the new call would fail to dispatch on replay. The
registration is added next to
`saveBuiltinModule`, mirroring the interface's append order.
- **`tools/slang-unit-test/unit-test-vtable-stability.cpp` (vtable
guard).** The probe must implement
every `IGlobalSession` method (pure virtuals), so the new method is
added as slot 32 (after
`saveBuiltinModule` = 31) and asserted. This is the regression guard
proving the new method is the
*last* slot — i.e. truly appended, not inserted.
- **`tools/slang-unit-test/unit-test-downstream-compiler-version.cpp`
(test).** A pure metadata
query, so it runs without a GPU. It first asserts the invalid inputs —
`SLANG_PASS_THROUGH_NONE`
(with both real and null out-params, so null handling runs on every CI
runner) and an out-of-range
`SlangPassThrough(SLANG_PASS_THROUGH_COUNT_OF)` (which exercises the
boundary guard) — all →
`SLANG_E_NOT_FOUND`, before any availability probe so they cannot
perturb the lazy cache. It covers
the documented "loaded-but-versionless ⇒ `SLANG_OK` with `(0,0)`" clause
via
`SLANG_PASS_THROUGH_GLSLANG` (bundled, GPU-free, reports no numeric
version) so a future regression
conflating that with `SLANG_E_NOT_FOUND` would fail. Then, gated on
`checkPassThroughSupport(SLANG_PASS_THROUGH_NVRTC)` exactly as the
existing CUDA tests do, it
asserts `SLANG_OK` with `major > 0` (and null out-params tolerated) when
NVRTC is available, or
`SLANG_E_NOT_FOUND` when it is not. CI runners with NVRTC exercise the
populated path.
## CLI option `-<compiler>-version` (added per review request)
@jkwak-work asked to surface this new API through a slangc command-line
query option mirroring the
existing `-<compiler>-path`, plus slang-test coverage. That is the final
commit on this branch.
**Behavior.** `-<compiler>-version` is a print-and-continue query flag
with the same
`-<compiler>-...` shape as `-<compiler>-path`: it consumes no value,
prints
`<compiler> version: <major>.<minor>` for the downstream compiler Slang
would actually load for that
pass-through, then lets compilation continue (exactly like `-version`).
It honors `-<compiler>-path`
and the standard search order because it routes through the same
`getDownstreamCompilerVersion`
funnel. When the toolchain cannot be located it prints `<compiler>
version: not found` and still
exits 0, so the `<compiler> version:` prefix is stable across machines.
Examples:
```
slangc -dxc-version # -> "dxc version: 1.8" (or "dxc version: not found")
slangc -glslang-version # -> "glslang version: 0.0" (loaded, reports no numeric version)
slangc -nvrtc-version -dxc-version # prints both, in command-line order; consumes no value
```
**Why this layer.** The parser already owns the symmetric
`-<compiler>-path` registration/handler;
the new option reuses the same `NameValueUtil::getNames(...
getCompilerInfos())` name expansion and
the same `lastIndexOf('-')` name recovery, and the only new sink call is
`diagnoseRaw(Note, ...)`
(verbatim text, no error-count increment, exit stays 0). No new public
surface beyond the appended
`CompilerOptionName::CompilerVersion = 153` enumerator (CLI-only; never
stored on an option set).
| File | Change |
|------|--------|
| `include/slang.h` | Append `CompilerOptionName::CompilerVersion = 153`
before `CountOf` (ABI-additive; CLI-only). |
| `source/slang/slang-options.cpp` | Register `-<compiler>-version`
(mirrors the `-<compiler>-path` block) and add the print-and-continue
handler. |
| `tests/downstream/downstream-compiler-version.slang` (new) |
Comment-only `//TEST:SIMPLE(filecheck=...)` cases for
dxc/fxc/glslang/nvrtc, the hyphenated `spirv-dis` name path, and a
two-flag invocation — matching the stable `<compiler> version:` prefix
so they are deterministic on any runner (GPU-free). |
| `docs/command-line-slangc-reference.md` | Regenerated from `slangc
-help-style markdown -h` to include the new option. |
Closes #11552.
---------
Co-authored-by: nv-slang-bot[bot] <274397474+nv-slang-bot[bot]@users.noreply.github.com>
Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
1 parent 168c59b commit 0d8e757
10 files changed
Lines changed: 290 additions & 1 deletion
File tree
- docs
- include
- source
- slang-record-replay
- proxy
- slang
- tests/downstream
- tools/slang-unit-test
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
699 | 699 | | |
700 | 700 | | |
701 | 701 | | |
| 702 | + | |
| 703 | + | |
| 704 | + | |
| 705 | + | |
| 706 | + | |
| 707 | + | |
| 708 | + | |
| 709 | + | |
| 710 | + | |
| 711 | + | |
702 | 712 | | |
703 | 713 | | |
704 | 714 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
1204 | 1204 | | |
1205 | 1205 | | |
1206 | 1206 | | |
| 1207 | + | |
| 1208 | + | |
| 1209 | + | |
| 1210 | + | |
| 1211 | + | |
| 1212 | + | |
1207 | 1213 | | |
1208 | 1214 | | |
1209 | 1215 | | |
| |||
4192 | 4198 | | |
4193 | 4199 | | |
4194 | 4200 | | |
| 4201 | + | |
| 4202 | + | |
| 4203 | + | |
| 4204 | + | |
| 4205 | + | |
| 4206 | + | |
| 4207 | + | |
| 4208 | + | |
| 4209 | + | |
| 4210 | + | |
| 4211 | + | |
| 4212 | + | |
| 4213 | + | |
| 4214 | + | |
| 4215 | + | |
| 4216 | + | |
| 4217 | + | |
| 4218 | + | |
| 4219 | + | |
| 4220 | + | |
| 4221 | + | |
| 4222 | + | |
| 4223 | + | |
4195 | 4224 | | |
4196 | 4225 | | |
4197 | 4226 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
257 | 257 | | |
258 | 258 | | |
259 | 259 | | |
| 260 | + | |
| 261 | + | |
| 262 | + | |
| 263 | + | |
| 264 | + | |
| 265 | + | |
| 266 | + | |
| 267 | + | |
| 268 | + | |
| 269 | + | |
| 270 | + | |
| 271 | + | |
| 272 | + | |
| 273 | + | |
| 274 | + | |
| 275 | + | |
| 276 | + | |
| 277 | + | |
260 | 278 | | |
261 | 279 | | |
262 | 280 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
124 | 124 | | |
125 | 125 | | |
126 | 126 | | |
| 127 | + | |
127 | 128 | | |
128 | 129 | | |
129 | 130 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
274 | 274 | | |
275 | 275 | | |
276 | 276 | | |
| 277 | + | |
| 278 | + | |
| 279 | + | |
| 280 | + | |
| 281 | + | |
| 282 | + | |
| 283 | + | |
| 284 | + | |
| 285 | + | |
| 286 | + | |
| 287 | + | |
| 288 | + | |
| 289 | + | |
| 290 | + | |
| 291 | + | |
| 292 | + | |
| 293 | + | |
| 294 | + | |
| 295 | + | |
| 296 | + | |
| 297 | + | |
| 298 | + | |
| 299 | + | |
| 300 | + | |
| 301 | + | |
| 302 | + | |
| 303 | + | |
| 304 | + | |
| 305 | + | |
| 306 | + | |
| 307 | + | |
| 308 | + | |
277 | 309 | | |
278 | 310 | | |
279 | 311 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
154 | 154 | | |
155 | 155 | | |
156 | 156 | | |
| 157 | + | |
| 158 | + | |
| 159 | + | |
| 160 | + | |
157 | 161 | | |
158 | 162 | | |
159 | 163 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
981 | 981 | | |
982 | 982 | | |
983 | 983 | | |
| 984 | + | |
| 985 | + | |
| 986 | + | |
| 987 | + | |
| 988 | + | |
| 989 | + | |
| 990 | + | |
| 991 | + | |
| 992 | + | |
| 993 | + | |
| 994 | + | |
| 995 | + | |
| 996 | + | |
| 997 | + | |
| 998 | + | |
| 999 | + | |
| 1000 | + | |
| 1001 | + | |
| 1002 | + | |
| 1003 | + | |
| 1004 | + | |
| 1005 | + | |
984 | 1006 | | |
985 | 1007 | | |
986 | 1008 | | |
| |||
3639 | 3661 | | |
3640 | 3662 | | |
3641 | 3663 | | |
| 3664 | + | |
| 3665 | + | |
| 3666 | + | |
| 3667 | + | |
| 3668 | + | |
| 3669 | + | |
| 3670 | + | |
| 3671 | + | |
| 3672 | + | |
| 3673 | + | |
| 3674 | + | |
| 3675 | + | |
| 3676 | + | |
| 3677 | + | |
| 3678 | + | |
| 3679 | + | |
| 3680 | + | |
| 3681 | + | |
| 3682 | + | |
| 3683 | + | |
| 3684 | + | |
| 3685 | + | |
| 3686 | + | |
| 3687 | + | |
| 3688 | + | |
| 3689 | + | |
| 3690 | + | |
| 3691 | + | |
| 3692 | + | |
| 3693 | + | |
| 3694 | + | |
| 3695 | + | |
| 3696 | + | |
| 3697 | + | |
| 3698 | + | |
| 3699 | + | |
| 3700 | + | |
| 3701 | + | |
| 3702 | + | |
| 3703 | + | |
| 3704 | + | |
| 3705 | + | |
| 3706 | + | |
| 3707 | + | |
| 3708 | + | |
| 3709 | + | |
| 3710 | + | |
| 3711 | + | |
3642 | 3712 | | |
3643 | 3713 | | |
3644 | 3714 | | |
| |||
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
| 1 | + | |
| 2 | + | |
| 3 | + | |
| 4 | + | |
| 5 | + | |
| 6 | + | |
| 7 | + | |
| 8 | + | |
| 9 | + | |
| 10 | + | |
| 11 | + | |
| 12 | + | |
| 13 | + | |
| 14 | + | |
| 15 | + | |
| 16 | + | |
| 17 | + | |
| 18 | + | |
| 19 | + | |
| 20 | + | |
| 21 | + | |
| 22 | + | |
| 23 | + | |
| 24 | + | |
| 25 | + | |
| 26 | + | |
Lines changed: 91 additions & 0 deletions
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
| 1 | + | |
| 2 | + | |
| 3 | + | |
| 4 | + | |
| 5 | + | |
| 6 | + | |
| 7 | + | |
| 8 | + | |
| 9 | + | |
| 10 | + | |
| 11 | + | |
| 12 | + | |
| 13 | + | |
| 14 | + | |
| 15 | + | |
| 16 | + | |
| 17 | + | |
| 18 | + | |
| 19 | + | |
| 20 | + | |
| 21 | + | |
| 22 | + | |
| 23 | + | |
| 24 | + | |
| 25 | + | |
| 26 | + | |
| 27 | + | |
| 28 | + | |
| 29 | + | |
| 30 | + | |
| 31 | + | |
| 32 | + | |
| 33 | + | |
| 34 | + | |
| 35 | + | |
| 36 | + | |
| 37 | + | |
| 38 | + | |
| 39 | + | |
| 40 | + | |
| 41 | + | |
| 42 | + | |
| 43 | + | |
| 44 | + | |
| 45 | + | |
| 46 | + | |
| 47 | + | |
| 48 | + | |
| 49 | + | |
| 50 | + | |
| 51 | + | |
| 52 | + | |
| 53 | + | |
| 54 | + | |
| 55 | + | |
| 56 | + | |
| 57 | + | |
| 58 | + | |
| 59 | + | |
| 60 | + | |
| 61 | + | |
| 62 | + | |
| 63 | + | |
| 64 | + | |
| 65 | + | |
| 66 | + | |
| 67 | + | |
| 68 | + | |
| 69 | + | |
| 70 | + | |
| 71 | + | |
| 72 | + | |
| 73 | + | |
| 74 | + | |
| 75 | + | |
| 76 | + | |
| 77 | + | |
| 78 | + | |
| 79 | + | |
| 80 | + | |
| 81 | + | |
| 82 | + | |
| 83 | + | |
| 84 | + | |
| 85 | + | |
| 86 | + | |
| 87 | + | |
| 88 | + | |
| 89 | + | |
| 90 | + | |
| 91 | + | |
| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
665 | 665 | | |
666 | 666 | | |
667 | 667 | | |
668 | | - | |
| 668 | + | |
669 | 669 | | |
670 | 670 | | |
671 | 671 | | |
| |||
841 | 841 | | |
842 | 842 | | |
843 | 843 | | |
| 844 | + | |
| 845 | + | |
| 846 | + | |
| 847 | + | |
| 848 | + | |
| 849 | + | |
844 | 850 | | |
845 | 851 | | |
846 | 852 | | |
| |||
862 | 868 | | |
863 | 869 | | |
864 | 870 | | |
| 871 | + | |
| 872 | + | |
865 | 873 | | |
866 | 874 | | |
867 | 875 | | |
| |||
0 commit comments