| Age | Commit message (Collapse) | Author |
|---|
|
* Perserve specialization cache in IR for specialization pass.
* Fix compile error.
* Fix.
* Fix.
* Fix test case.
* Fix.
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* Lower throwing COM interface method.
* Fix.
* Fix warnings.
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* Use TerminatedUnownedStringSlice for literals in output C++.
* Remove Escape/Unescape functions used in slang-token-reader.cpp
Add target type of 'host-cpp' etc to map to the target types.
* Fix some corner cases around string encoding.
* Added unit test for string escaping.
Fixed some assorted escaping bugs.
* Updated test output.
* Added decode test.
* Stop using hex output, to get around 'greedy' aspect. Use octal instead.
* Added HostHostCallable
Small changes to use ArtifactDesc/Info instead of large switches.
* Fix C++ emit to handle arbitrary function export.
* Add options handling for callable without an output being specified.
* Can compile with COM interface. Added example using com interface.
* Use the IR Ptr type instead of hack in C++ emit for interfaces.
* Fix issue with outputting the COM call when ptr is used.
* Fix crash issue on compilation failure.
|
|
* New language feature: basic error handling.
* Fix.
* Fix `tryCall` encoding according to code review.
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* Remove the need for LivenessLocation.
* Use LivenessMode.
* Fix some comments.
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
See https://github.com/shader-slang/slang/issues/2213
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* Refactor how prelude output works in emit.
* Small improvement to emit output.
* Move around comment on target specific language directives based on review.
Co-authored-by: Theresa Foley <10618364+tangent-vector@users.noreply.github.com>
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* Refactor Liveness pass, such that locations can be found independently of setting up ranges.
* Refactor around different stages of liveness span analysis.
* WIP Take into account PHI temporaries in liveness tracking.
* WIP First pass of PHI liveness refactor.
* Add BlockIndex.
* WIP Refactor phi liveness around inst runs.
* More improvements around liveness tracking.
* Bug fixes.
Special handling to not add multiple ends, at starts of blocks and after accesses.
* Fix test output.
* Use IRInsertLoc to track insertion point.
* Liveness markers don't have side effects.
* Fix typo in liveness test.
* Small improvements around setting SuccessorResult.
* Fix memory issue around reallocation and RAIIStackArray.
Update test output.
* Update test output for liveness.slang.
* Fix typo in SuccessorResult blockIndex.
* Small tidy up.
* Handle the root start block, correctly scoping the run.
* Split BlockInfo into 'Root' and 'Function'.
Store successors as BlockIndices.
* Tidy up around liveness tracking.
* Add head/tail support to ArrayViews.
Use Count where appropriate.
Use head/tail in liveness impl.
|
|
Co-authored-by: Yong He <yhe@nvidia.com>
Co-authored-by: Theresa Foley <10618364+tangent-vector@users.noreply.github.com>
|
|
* Use IR pass to eliminate phi nodes
"Phi nodes" are one of the key contrivances that makes SSA (Static
Single Assignment) form work. Because SSA is so great for compiler
IRs, we kind of need to deal with phi nodes, but they also get in
the way because they don't have a direct analog in most lower-level
machine ISAs or execution models, nor in most of the high-level
languages a transpiler wants to emit. As a result a compiler like
ours needs to be able to eliminate the phi nodes from a program as
part of generating output code.
(For any clever people noting that SPIR-V supports phi nodes
directly: yes, it does. It doesn't need to and it probably *shouldn't*.
Anybody involved in the decision-making knows my reasoning, and
anybody else should feel free to ask me if they want the lecture.
Anyway...)
The basic idea of elimiating phi nodes is simple enough. We replace
each phi node with a temporary variable. Uses of the phi use values
loaded from the temporary. The operation of the phi itself
(assigning a value based on the branch taken) amounts to an assignment
into the temporary.
Previously, the Slang compiler dealt with phi nodes very late in
the process of generating code: in the middle of emitting strings
of source code in a high-level language like HLSL or GLSL. Doing the
work that late in compilation has two big drawbacks:
1. Our ability to emit clean and/or optimal code is limited because
we may not be able to make certain changes to the IR, or because we
cannot make use of additional information like a dominator tree that
might be available at other points in compilation.
2. Any other IR passes that relate to temporary variables won't be
able to see the variables that we generate for phi nodes. This could
raise issues with correctness (e.g., if we want to compute live-range
information for *all* temporary variables), or performance (we have no
way to run additional IR optimization passes after phis are eliminated).
This change addresses these problems by making the elimination of
phi nodes an explicit IR pass. Additional optimizations can easily be
run after this pass (although we'd need to be careful not to run
passes that could end up introducing new phis). The pass makes use
of the information available to it to try to produce code that will
emit to "clean" HLSL/GLSL.
The core of the pass is in `slang-ir-eliminate-phis.cpp`, and is
heavily commented, so I won't describe the approach in detail here.
There are two related issues that came up, though:
First, it turned out that our emit logic for local variables (`IRVar`
instructions) wasn't using the function we'd defined named `emitVar()`.
One worrying consequence of that oversight was that the `precise`
modifier would impact generated HLSL/GLSL for variables that turned
into SSA values (including phi nodes), but *not* for local variables
that had not been SSA'd (or that had been SSA'd and then de-SSA'd).
This change also fixes that bug; it is unclear how widespread the
impact of the original issue might be.
Second, generating explicit IR temporaries for phi nodes exposed a
pre-existing bug in the `slang-ir-restructure-scoping` pass. That pass
basically detects cases where we have an instruction `I` with a use
`U` such that the use follows the rules of SSA form ("def dominates
use," meaning `I` dominations `U`), but does not follow the more
restrictive scoping rules of high-level-language output (where a value
computed "inside" a loop is not automatically visible to code outside
the loop just because it dominates that code). That pass did not
correctly account for the case where `I` was a temporary variable.
It seems that case could not arise before now because we didn't have
any passes that would move `var`, `load`, or `store` operations out
of the basic block they started in. The fix for that pass was relatively
simple, and will make the whole thing more robust in case we add more
aggressive optimizations later.
* fixup: expected test output
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* WIP tracking liveness.
* Skeleton around adding liveness instructions.
* Calling into liveness tracking logic.
Adds live start to var insts.
* Liveness macros have initial output.
* Looking at different initialization scenarios.
* Some discussion around liveness.
* WIP for working out liveness end.
* WIP Updated liveness using use lists.
* Is now adding liveness information
* Some small fixes.
* WIP around liveness.
* Seems to output liveness correctly for current scenario.
* Tidy up liveness code.
* Update comment arounds liveness to current status.
* Small fixes to liveness test.
* Add support for call in liveness analysis.
* Improve liveness example with array access.
* Small updates to comments.
* Disable liveness test because inconsistencies with output on CI system.
* First pass support for GLSL SPIR-V liveness support.
* Add the SPIRVOpDecoration.
* Fix signature for OpLivenessStop.
* Simplified by having a Kind type.
* Fix some issues brought up in PR.
* Rename liveness instructions.
* Merge with var-lifetime.
Small improvements.
* Improvements to the documentation/naming in GLSL liveness pass.
Add comment around possible improvements to the liveness pass.
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* WIP tracking liveness.
* Skeleton around adding liveness instructions.
* Calling into liveness tracking logic.
Adds live start to var insts.
* Liveness macros have initial output.
* Looking at different initialization scenarios.
* Some discussion around liveness.
* WIP for working out liveness end.
* WIP Updated liveness using use lists.
* Is now adding liveness information
* Some small fixes.
* WIP around liveness.
* Seems to output liveness correctly for current scenario.
* Tidy up liveness code.
* Update comment arounds liveness to current status.
* Small fixes to liveness test.
* Add support for call in liveness analysis.
* Improve liveness example with array access.
* Small updates to comments.
* Disable liveness test because inconsistencies with output on CI system.
* Fix some issues brought up in PR.
* Rename liveness instructions.
|
|
* Support `[DllImport]`
* Fix.
* Fix.
* Fix array type emit in cpp.
* Fix.
* Fix.
* Fix
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
An earlier refactoring pass over the compiler codebase split the
type that had been called `CompileRequest` into three distinct
pieces:
* `FrontEndCompileRequest` which was supposed to own state and
options related to running the compiler front end and producing
IR + reflection (e.g., what translation units and source
files/strings are included).
* `BackEndCompileRequest` which was supposed to own state and options
related to running the compiler back end to translate the IR
for a `ComponentType` (program) into output code. (Note that the
`BackEndCompileRequest` was conceived of as orthogonal to the
`TargetRequest`s, which store per-target and target-specific
options.)
* `EndToEndCompileRequest` which was an umbrella object that owns
separate front-end and back-end requests, plus any state that is
only relevant when doing a true end-to-end compile (such as the
kinds of compiles initiated with `slangc`). As originally conceived,
the only state that this type was supposed to own was stuff related
to "pass-through" compilation, as well as state related to writing
of generated code to output files.
That refactoring work was very useful at the time, because it allowed
us to "scrub" the back end compilation steps to remove all
dependencies on front-end and AST state (this was important for our
goals of enabling linking and codegen from serialized Slang IR).
At this point, however, it is clear that the hierarchy that was built
up serves very little purpose:
* The `BackEndCompileRequest` type is only used in two places:
* As part of an `EndToEndCompileRequest`, where the settings on
the `BackEndCompileRequest` can be configured, but only through
the `EndToEndCompileRequest`
* As part of on-demand code generation through the `IComponentType`
APIs. In this case, the settings stored on the
`BackEndCompileRequest` are not accessible to the application
at all, and will always use their default values, so that
instantiating a "request" object doesn't really make any sense.
* The `FrontEndCompileRequest` type has a similar situation:
* Front-end compilation as part of an `EndToEndCompileRequest`
supports user configuration of `FrontEndCompileRequest` settings,
but only through the `EndToEndCompileRequest`
* Front-end compilation triggered by an `import` or a `loadModule()`
call does not support user configuration of settings at all. It
will always derive all relevant settings from thsoe on the
session ("linkage").
In addition, subsequent changes have been made to the compiler that
show a bit of a "code smell" and/or forward-looking worries for this
decomposition:
* In some cases we've had to add the same setting to multiple types
in the breakdown (front-end, back-end, end-to-end, linkage, target,
etc.) which makes it harder for us to validate that all the possible
mixtures of state work correctly.
* Related to the above, in some cases we have manual logic that copies
state from one of the objects in the breakdown to another, in order
to ensure that the user's intention is actually followed.
* As a forward-looking concern, it seems that developers have sometimes
added new configuration options and state to places that don't really
make sense according to the rationale of the original decomposition
(e.g., we probably don't want to have a lot of state that is only
available via end-to-end requests, given that the API structure is
meant to push users *away* from end-to-end compiles).
As a result of all of the above, I've been planning a large refactor
with the following big-picture goals:
* Eliminate `BackEndCompileRequest`
* Move all relevant state/options from the back-end request to
the end-to-end request, since that is the only place they could
be set anyway.
* Introduce a transient "context" type to be used for the duration
of code generation that serves the main functions that back-end
requests really served in the codebase
* Make `EndToEndCompileRequest` be a subclass of
`FrontEndCompileRequest`
* Consider addding a transient "context" type for front-end
compiles that can be used in `import`-like cases rather than
needing a full front-end request object. If this works, then
eliminate `FrontEndCompileRequest` and be back to world with
just a single `CompileRequest` type
* Move *all* compiler configuration options to a distinct type (named
something like `CompilerConfig` or `CompilerOptions` or whatever)
which stores setting as key-value pairs, and has a notion of
"inheritance" such that one configuration can extend or build on top
of another. Make all the relevant types use this catch-all structure
instead of redundantly storing flags in many places.
This change deals with the first of those bullets: removeal of
`BackEndCompileRequest`. The addition of the `CodeGenContext` type is
perhaps an unncessary additional step, but making that change helps
clean up a bunch of the code related to per-target code generation,
so I think it is the right choice.
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
|
|
`ImageSubscript` for GLSL (#2146)
|
|
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* Various gfx fixes.
* Fix test case.
* Fix crash.
* Trigger build
* Trigger build 2
* Fix vulkan unit tests.
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* WIP control of dump options.
* Removed SourceManager for IRDumpOptions
* Arm aarch64 debug connection timeout - as CI timed out.
|
|
* `reinterpret` and 16-bit value packing.
* Update `half-texture` cross-compile test reference result.
* Revert inadvertent reformatting of slang-ir-inst-defs.h
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* Further implementation of SPIRV direct emit.
This change implements:
- Struct, Vector, Matrix and Unsized Array types.
- Basic arithmetic opcodes, vector construct, swizzle etc.
- getElementPtr, getElement, fieldAddress, extractField.
- SPIRV target intrinsics with SPIRV asm code in stdlib.
- RWStructuredBuffer and StructuredBuffer.
- Pointer storage class propagation.
- Control flow.
* Fix.
|
|
* Fix a few issues around opaque types as outputs
Slang and HLSL support opaque types (textures, buffers, samplers, etc.) as members of `struct`s, mutable local variables, function results, and `out`/`inout` parameters. GLSL and SPIR-V do not.
In order to translate Slang code over to GLSL/SPIR-V we use a variety of passes that seek to eliminate all of the above use cases and produce code that only uses opaque types in the limited ways that GLSL/SPIR-V allow. This change relates to the passes that deal with function results and `out`/`inout` parameters.
There are two basic changes here:
1. The `specializeResourceOutputs` pass was only dealing with resource (texture/buffer) types. This change updates it to process sampler types as well.
2. The sequencing of the passes made it possible that an opaque-typed local variable might be left around after `specializeResourceOutputs`, which would mean the code is still invalid for GLSL/SPIR-V. This change adds an additional SSA-formation pass which would eliminate any opaque-type local variables whose lifetimes were made simple enough by the optimizations.
Together these changes fix a problem-case user shader that was failing to compile for Vulkan.
* Update slang-emit.cpp
Fix typo 'reuslt'
* Update slang-emit.cpp
Comment change to re-trigger CI build.
Co-authored-by: jsmall-nvidia <jsmall@nvidia.com>
|
|
points (#1917)
* optix SBT record data can now be accessed using uniform parameters on ray tracing entry points
* Update slang-emit.cpp
|
|
* Work to mitigate SPIR-V bloat
SPIR-V is not an especially compact format, but some patterns in how Slang generates code and then runs it through `spirv-opt` lead to many redundant field-by-field copy operations being emitted. This change attempts to address some of the resulting bloat from the Slang side of things.
Note: experimentation shows that the bloat is less pronounced when running either *no* SPIR-V optimizations or *full* SPIR-V optimizations, so it is also likely that the bloat should be addressed by changing which `spirv-opt` passes the Slang compiler runs in default (`-O1`) builds. Such changes should come as a distinct pull request.
This change primarily does two things:
First, the code generation strategy for passing arguments to `out` and `inout` parameters has been changed. In the past, the compiler would *always* copy the argument value into a temporary, then pass the address of the temporary, and then write back the value after the call. The new code generation strategy attempts to identify when an argument value already has a simple address in memory and passes that address directly when possible. This eliminates many copy operations that occur before/after calls to functions with `out`/`inout` parameters.
Second, we introduce an IR optimization pass that detects call sites where the entire contents of a buffer (usually a constant buffer) is being passed to a callee function, such that many bytes are loaded and then passed even if only very few are used in the callee. The pass moves the load operations from the caller to a specialized version of the the callee where possible (e.g., when the constant buffer in question is a global shader parameter). Doing this eliminates another major category of copies.
Notes:
* The IR lowering logic is complicated by the fact that several kinds of l-values (values that are usable as the desitnation of assignment, or for `out`/`inout` arguments) are not actually addressable. An easy example is a non-contiguous swizzle like `v.xwz` on a `float4`, where the value occupies 12 bytes, but not 12 consecutive bytes with a single address. There are many more corner cases like that and the IR lowering pass carries a lot of complexity to deal with them. A more systematic overhaul is due some time soon.
* The IR representation of `out` and `inout` parameters deserves some careful scrutiny when making these kinds of changes. The official semantics of `inout` in HLSL has been "copy in copy out" (and `out` is just "copy out") which is observably different from any solution that passes in the address of an l-value directly. By making this change we are saying that Slang's semantics are not precisely those of legacy HLSL, and that our semantics for `inout` parameters are closer to those of `inout` in Swift or of a mutable borrow in Rust. In the Swift case the implementation can freely pass the underlying storage of an l-value or the address of a temporary, and valid programs may not observe the different. It is thus illegal to observe the value in a storage local while a mutation to that location is "in flight." All of this is way more detailed and technical than 99% of Slang users will ever care about, but importantly it gives us semantic cover to eliminate these copies in the IR, and also to emit output C++ code that implements `out` and `inout` as by-reference parameter passing.
* There was an exsting generic pass for specializing functions based on call sites that uses a "template method" style of pattern to customize its behavior. That pass needed to be generalized to handle this use case because it had previously operated on the assumption that the "desire" to specialize a callee function must be driven by the parameter declarations of that function, and not on the argument values passed in. The code has been slightly refactored to allow the policy for specialization to consider both parameters and arguments.
* Unsurprisingly, a bunch of the GLSL (and thus SPIR-V) generated has changed with this work, so several baseline `.slang.glsl` files needed to be updated.
* This change is incomplete in that it does not address broader cases of buffer loads, including both partial loads from constant buffers (just loading one field, but a field that uses a "large" structure type), and loads from multi-element buffers (a lot from a structured buffer where the element type is "large"). The main question in each of those cases is how to define how "large" a structure needs to be before we decide to try and sink loads into callee functions like this. In the worst case, sinking loads in this way may actually create *more* memory traffic (because the same values get loaded in multiple callee functions).
* fixup: run premake
* fixup: typo
|
|
* Various Fixes to gfx, reflection and emit.
- Fix GLSL emit to properly output `*bitsTo*` functions for `IRBitCast` insts.
- Add line directive mode setting for `ISession`.
- Extend `TypeLayout::getElementStride` to handle `VectorType` case.
- Fix `IDevice::readBufferResource` 's D3D12 implementation to copy only the requested bytes out.
- Fix `render-test` to use the `ISession` from `gfx` instead of creating its own `ISession` to make sure `gfx` and `render-test` agree on WitnessTable and RTTI IDs.
- Extend `render-test` to support filling vector and matrix values in the new `set x = ...` TEST_INPUT syntax.
- Add a `dynamic-dispatch-15` test case to make sure packing / unpacking works correctly across all targets, and to make sure render-test's RTTI/WitnessTable ID filling logic is correct for non-trivial cases.
* Remove default-major test
* Fix cyclic reference in `ExtendedTypeLayout`.
* Move `lineDirectiveMode` setting to `TargetDesc`.
Add `structureSize` to `TargetDesc` and `SessionDesc` for future binary compatibility.
* Cleanup.
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* Fix issue with with SLANG_ENABLE_GLSLANG_SUPPORT
* Update expected output from glslang-error.glsl
* Fix bug in glsl dissassembly.
* Make ExtensionTracker available even if source is not emitted.
* Only explicitly set extension tracker based on capability bits, if we are in pass through.
* Small simplification of invoke sourceEmit.
|
|
During lowering from AST to IR, the Slang compiler translates code that uses `struct` inheritance:
```hlsl
struct Base { int a; }
struct Derived : Base {}
```
into code where the inheritance relationship is "witnessed" by a simple field:
```hlsl
struct Base { int a; }
struct Derived { Base __anonymous_field__; }
```
The underlying bug here is that the `__anonymous_field__` that the compiler generated during IR lowering was not being given any linkage decorations (no mangled name). As a result, if multiple separately-compiled modules all access that field they could disagree on its identity as an IR instruction. This could lead to output code being generated where the declaration of `__anonymous_field__` uses one IR instruction, but accesses use another.
This change includes a fix for the issue, and a test that serves as a reproducer for the original problem.
|
|
Co-authored-by: T. Foley <tfoleyNV@users.noreply.github.com>
|
|
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* Split out compiler-core initially with just slang-source-loc.cpp
* More lexer, name, token to compiler-core.
* Split Lexer and Core diagnostics.
* Move slang-file-system to core.
* Add slang-file-system to core.
* More DownstreamCompiler into compiler-core
* Fix typo.
* Add compiler-core to bootstrap proj.
* Small fixes to premake
* For linux try with compiler-core
* Remove compiler-core from examples.
* Added NameConventionUtil to compiler-core
* Add global function to CharUtil to *hopefully* avoid linking issue.
* Hack to make linkage of CharUtil work on linux.
|
|
* Add command-line control over SPIR-V version
By default the Slang compiler policy is usually to produce output with the fewest dependencies possible. If input code can be encoded as SPIR-V 1.0, that is what we will use by default. The catch here is that in some cases later SPIR-V versions introduced improvements to the encoding that can affect performance (e.g., around large global arrays of constants), so that a user might explicitly want to require a newer SPIR-V version (restricting the driver versions their code can work on) in the hopes of seeing better performance.
This change uses the system of capabilities that was previously introduced so that an option like `-profile glsl_450+spirv_1_5` can be used to explicitly request a specific SPIR-V version. Consistent with the existing implementation, the requested version will be taken as a minimum, and the final version might be higher based on other requirements (e.g., use of intrinsic functions that require a higher version).
The test case included here is a little iffy in terms of long-term maintanenace. It relies on having both a `.slang` file and a `.glsl` file that we compile with the same options and then compare the SPIR-V, but that means there is no direct testing that the output SPIR-V actually uses the necessary version. If we break the inference of SPIR-V versions for both the regular and pass-through paths at once, this test won't flag the problem. A better test is probably needed soon.
This change *only* adds support for controlling the SPIR-V version via capabilities specified via the command line or API. It would be nice to a future change to allow something like `[require(spirv_1_5)]` to be added to an entry point function to allow the user to embed their expectation/requirement into the source code.
* fixup: clang warning
|
|
* #include an absolute path didn't work - because paths were taken to always be relative.
* WIP: First pass in supporting output of line error information.
* Add support for lexing to better be able to indicate SourceLocation information.
* Fix lexer usage in DiagnosticSink in C++ extractor.
* Update diagnostics tests to have line location info.
* Fixed test expected output that now have source location information in them.
* Better handling of tab.
* Fix test expected results for tabbing change.
* DiagnosticLexer -> DiagnosticSink::SourceLocationLexer
Added line continuation tests.
* Fix typo.
* Added String::appendRepeatedChar
* Change to rerun tests.
* Added source locations to IR dumping.
* Output column for IR dump source loc.
* Add support for closing brace location to AST.
Use closing brace location in lowering when adding return void.
* Set the source location through SourceLoc - simplifies identifying if current loc is valid.
* Copy terminator sloc.
* Test for improved #line handling.
* Made writer the last parameter for dumpIR.
Small improvements to comments.
* Disable sloc output on dump IR by default.
* Fix issue with #line and inlining.
* Fix for output with improved #line output.
* Small comment change - mainly to kick off TC build.
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Support `bit_cast` between complex types.
* Fix vs project file
* Fix clang build error
* fix
* fix
* Fix
* FIx
* Fix
* Fix
* Fix
* Fix
* Fix linux compile error
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Use "capability" system to select VKRT extension
Slang currently supports translation of ray tracing shader code to Vulkan GLSL code that uses the `GL_NV_ray_tracing` extension. A multi-vendor equivalent of that extension has been released as `GL_EXT_ray_tracing` and we want Slang to support that extension as well.
At the simplest, making the change from one extension to the other is just a matter of changing a few strings, since it does not appear that anything of significance was changed at the GLSL level (or even in SPIR-V). Where this gets trickier is when we have users who want us to support *both* extensions, and to be able to switch between them.
The solution we've implemented here more or less amounts to:
* If you don't tell the compiler which extension to use, it will default to `GL_EXT_ray_tracing` (the newer multi-vendor one).
* If you explicitly want the older extension, you can opt into it using the `-profile` option or via a new API for explicitly adding capabilities to your target.
Making that work required a few different kinds of changes:
* The options parsing and public API needed ways to add optional capabilities to a target.
* During GLSL code emit, we can check the capabilities that were added to the target to see if the `GL_NV_ray_tracing` extension was explicitly enabled and, if not, default to using the `GL_EXT_ray_tracing` names for things. This step is needed because some of the modifiers/attributes involved in the extension have to be handled explicitly in the code generator rather than implicitly as part of mapping intrinsic functions.
* We add two different translations to the relevant operatiosn in the stdlib, one marked with each of the extensions. If profile/capability-based overload resolution can be relied on to pick the right one, this should Just Work.
* Next, a bunch of work had to go into making capability-based overloading Just Work for the purposes of this change. There's been a nearly complete reworking of the implementation of `CapabilitySet` here to make it more suitable for our needs.
* The tests that were using ray tracing translation for Vulkan needed to be updated. For some of them I updated their baselines to use `GL_EXT_ray_tracing` so that they can test the new path. For others, I updated the command line for the test case so that it explicitly opts into using `GL_NV_ray_tracing`. The result is that we have some coverage of each extension. I would have liked to have each test run in both modes, but our pass-through glslang support doesn't support `-D` options, so I couldn't take that step easily.
This change does *not* add support for `GL_EXT_ray_query`, the extension that supports "DXR 1.1" style queries under Vulkan. Adding support for that extension should hopefully be a smaller step because it doesn't have the same multiple-extensions issue.
This change does *not* address a lot of possible avenues for improvement or cleanup around the capability system. It focuses only on those changes that are necessary to make the ray tracing feature work and leaves the rest for future work.
* fixup: infinite loop
* Comment-only change to retrigger TC build
|
|
* PR to fix issue #1638. This change introduces a diagnostic sink to the
emitModule function, and updates all associated calls to that function.
Additionally, this commit updates the heterogeneous hello world example
to not need the entry and stage flags for simplicity.
* Updated emit-cpp per suggested changes
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Add first steps toward a "capability" system
We already have cases in the stdlib where we mark declarations as being specific to certain targets, e.g.:
```
// My ordinary function to add two numbers.
// Works everywhere.
//
void myFunc(int a, int b) { return a + b; }
// On the "coolgpu" target, we can use a secret intrinsic
// that adds numbers even faster!
//
__specialized_for_target(coolgpu)
void myFunc(int a, int b) { return __secretIntrinsic(a, b); }
```
The existing logic for dealing with these modifiers (`__specialized_for_target` and `__target_intrinsic`) was almost entirely string-based. We would turn the chosen compilation target into a string, and then use that to try and search for the "best" definition of a function at a few steps:
* During IR linking, we always pick one definition of an `[import]`ed function, and that definition will be the one with the "best" target-specialization modifier (if any)
* During final code generation, we always look up the "best" target-intrinsic modifier, and use it as the template for the code we output.
This change preserves the basic flow there, but replaces the ad hoc string-based logic with something a bit more principled, in terms of a new `CapabilitySet` type.
A `CapabilitySet` represents a set of zero or more atomic features (here represented as `CapabilityAtom`s). What a `CapabilitySet` means depends on how and where it is used:
* A compilation target implies a `CapabilitySet` where the contents of the set are the features the target *supports*.
* A `CapabilitySet` attached to a declaration (or a modifier on that declaration) describes a set of feature that declaration *requires*.
The current implementation of `CapabilitySet` is wasteful and inefficient, but that is something we can iterate on over time.
In practice, most of the current code only ever uses capability sets that are either empty (because they represent a function with no specific requirements) or singleton (because they represent asingle atomic capability like "is a GLSL target," "is an HLSL target," etc.).
The main goal here was to put in the skeleton of a new system, including some of the features it might need down the line, and then to leave changes that eventually use the greater flexibility for later. Eventually, the capability system should encompass:
* Differences between shader model versions, GLSL versions, SPIR-V versions, etc. (currently tracked with other modifiers)
* Optional extensions, and functions that are made available only with certain extensions (currently tracked with other modifiers)
* Front-end checking that the call graph of a program doesn't violate any capability-requirements (e.g., having a GLSL+HLSL portable function call a GLSL-only subroutine)
* Hypothetically we can also try to fold stage-specific (vertex-only, fragment-only, etc.) functions into this system, but doing so would require more linker cleverness if we allow overloading on stages (since we might have to clone a caller if it calls through to a callee with multiple stage-specific versions)
One important complication that the system has to deal with just because of the "do what I mean" nature of the current compiler is that somethings a current Slang user might compile for target X and specify version N, but then use a function that actually requires version N+1 of that target. Currently the Slang compiler silently "upgrades" the version(s) used by user code in these cases, because it is often what users want in cross-compilation scenarios.
Dealing with the "silent upgrade" situation requires us to be a little careful and sometimes pick a "best" capability set that doesn't appear to be supported on our target. Refining that system and potentially getting rid of the "do what I mean" behavior over time could be a goal for future changes.
* fixup: handle case where value is incompatible during linking
|
|
Overview
========
Prior to this change, we had two different code generation strategies for interface/existential types in Slang, that didn't always play nicely together:
* The "legacy" static specialization approach could handle plugging in an arbitrary concrete type for an existential type parameter (including types with resources, etc.), but wouldn't work well with things like a `StructuredBuffer<>` of an interface type, and requires somewhat counter-intuitive layout rules to make work.
* The new dynamic dispatch approach produces simpler, more easily understood layouts by assuming that values of interface type can fit into a fixed number of bytes. The tradeoff there is that it cannot handle types that include resources (only POD types).
The goal of this change is to make it so that the two strategies can co-exist. In particular, in cases where a shader is amenable to both static specialization and dynamic dispatch, the type layouts should agree.
In order to make the type layouts agree, we:
* Declare that *all* values of existential type reserve storage according to the dynamic-dispatch rules (so 16 bytes for the RTTI and witness-table information, plus whatever bytes are needed to story "any value" of a conforming type).
* Then we modify the "legacy" layout rules so that if a value of concrete type can fit in the reserved "any value" space for a given interface, then it is laid out there exactly like the dynamic dispatch rules would do. Otherwise, we fall back to the previous legacy rules (since we don't need to agree with the dynamic-dispatch layout on types that can't be used with dynamic dispatch).
Details
=======
* Renamed `ExistentialBox` to `BoundInterfaceType` to better clarify how it relates to `BindExistentialsType`
* Unconditionally apply the `lowerGenerics` pass during emit, since it is now responsible for aspects of the lowering of existential types when specialization is used.
* Made IR type layout take the target into account, so that the layout of resource types can vary by target (e.g., being POD on some targets, and invalid on others)
* Cleaned up some issues around using global shader parameters as the "key" for their layout information in the global-scope layout (only comes up when there are global-scope `uniform` parameters)
* Made there be a default any-value size (16) instead of making it be an error to leave out. This was the simplest option; we could try to go back to having an error, but we'd need to only issue it if we are sure a type/interface is being used with dynamic dispatch, since static dispatch doesn't have to obey the restrictions.
* Changed lowering of existential types to tuples so that bound interfaces where the concrete type won't fit use a "pseudo-pointer" instead of an "any-value" to hold the payload
* Changed IR type legalization to handle the "pseudo-pointer" case and apply layout information from an interface type over to the payload part when static specialization was used.
* Changed some details of how witness tables were being lowered, so that we didn't have to create "proxy" witness tables for the constraints on associated types (just use the actual requirement entries we generate)
* Changed witness tables so that they know the subtype doing the conforming
* Added logic so that we don't generate pack/unpack logic and witness table wrapper functions for types that are incompatible with any-value/dynamic dispatch for a given interface.
* Changed the core AST-level type layout logic to use the dynamic-dispatch layout in case things fit, and the legacy static specialization case when things don't (while also reserving space for the dynamic-dispatch fields)
* Changed a bunch of test cases for static specialization to properly use the new layout (which introduces new buffers in some cases, and moves data around in others).
Future Work
===========
The experience of trying to reconcile our older way of handling interface-type specialization with our newer model (that supports dynamic dispatch) makes it clear that we really need to make similar changes to our handling of generic type parameters on entry points and at the global scope.
A future change should make it so that a global type parameter is lowered with a type layout similar to a value parameter of interface type, including the RTTI and witness-table pieces, and just leaving out the "any value" piece. A similar translation strategy should apply to entry-point generic parameters (mirroring how we lower generic functions for dynamic dispatch already), and value specialization parameters.
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* Use integer RTTI/witness handles in existential tuples.
* Fix clang error.
* Fix IR serialization to use 16bits for opcode.
* Undo accidental comment change.
* Use variable length encoding for opcode.
* Fix compile error.
* Fixing issues
* Fix code review issues.
|
|
In some cases, functionality is available as either a GLSL extension for Vulkan/SPIR-V, or through the NVAPI system for D3D. This situation creates complications because while GLSL extensions are generally all supported by the open-source glslang compiler (which we can bundle and ship), NVAPI operations are exposed through a specific header (`nvHLSLExtns.h`) that ships as part of the NVAPI SDK.
When a user wants to explicitly use NVAPI-provided operations in their shader code, there are no major complications for Slang; the user sets up their include paths, `#include`s the relevant header, calls functions in it, and lets Slang deal with the details of compilation.
The challenge for Slang arises when we want to provide a cross-platform interface in our standard library (e.g., the `RWByteAddressBuffer.InterlockedAddF32` method that was recently added) that uses either a GLSL extension (when compiling for Vulkan/SPIR-V) or an NVAPI (when compiling to DXBC or DXIL). In that case, the code *generated* by Slang now has a dependency on NVAPI, and we need to somehow emit a `#include` directive that pulls it in when invoking fxc or dxc. Because we do not (and seemingly cannot) bundle the NVAPI header with the compiler, we have to rely on ther user to have it available and to somehow communicate to Slang where it is.
Exposing portable routines that sometimes use NVAPI currently creates two main challenges:
1. The user is forced to interact with the "prelude" mechanism in the compiler, which allows the programmer to define code in a given target language that gets prepended to the Slang-generated code. While the prelude mechanism is powerful, it is also hard for users to integrate into their workflow, and our experience so far is that users want something that Just Works.
2. If the user writes code that uses some of our abstract operations that layer on NVAPI *and* they also want to use NVAPI explicitly, they end up with two copies of the NVAPI header (one included by the Slang front-end, and another included by the downstream fxc/dxc compiler). This puts the user in the situation of (a) having to ensure that they set the defines like `NV_SHADER_EXTN_SLOT` consistently both when invoking Slang and when adding their prelude, and (b) even if they do make the definitions consistent, they run into the problem that fxc/dxc complain about overlapping register bindings on the two copies of the `g_NvidiaExt` global shader paraemter that the NVAPI header declares.
This change attempts to resolve both issues by adding a lot of "do what I mean" logic to the compiler to try to ease things in the common case. In particular:
1. The user no longer needs to use the "prelude" mechanism when using NVAPI. The compiler now embeds a default prelude for HLSL output, which will `#include` the NVAPI header if and only if the generated code needs NVAPI access because of portable standard library routines that were used.
2. The user can mix-and-match explicit NVAPI use and stdlib functions that compile to use NVAPI. The register/space to be used by NVAPI when included via prelude is now set based on whatever the user set via the preprocessor so that it should automatically be consistent between both cases. Furthermore, the code we emit for the declaration of `g_NvidiaExt` when compiling explicit NVAPI use is set up to be conditional, so that it is skipped in the case where the prelude will pull in its own declaration of that parameter.
The way all this is achieved involves a lot of moving pieces:
* We now have an HLSL prelude, which mostly just serves to `#include "nvHLSLExtns.h"` in the case where NVAPI support is needed downstream.
* Standard library operations that require NVAPI for their implementation on HLSL include a new `[__requiresNVAPI]` attribute.
* The preprocessor has been extended so that after tokenizing an input file it looks up the NVAPI-relevant macros in the resulting environment, and if they are set it attached a modifier (`NVAPISlotModifier1) to the AST `ModuleDecl` that is based on their values. Logic is added to detect if multiple input files specify values for the macros in ways that conflict.
* The semantic checking step is extended so that it detects the "magic" NVAPI declarations (the `g_NvidiaExt` paramter and the `NvShaderExtnStruct` type that it uses) and attaches a modifier to them so that they can be identified as such in later steps.
* Parameter binding is extended to collect a list of the AST modifiers that reflect NVAPI binding, and to reserve the relevant register(s) so that ordinary user-defined parameters cannot conflict with them.
* IR lowering translates the three new AST modifiers related to NVAPI over to IR equivalents.
* IR linking is extended to make sure that it clones any `IRNVAPISlotDecoration`s attached to the input modules. The pass intentionally does not care where the modifiers came from; it just collects them all and leaves it to downstream code to sort out what they mean.
* Emit logic is extended to have a notion of "prelude directives" which are preprocessor directives that should come *before* the prelude in the generated code, because they can impact the way that the prelude compiles. This is done so that we don't have to introduce ad hoc logic for each downstream compiler to set any relevant `-D` flags (e.g., both fxc and dxc would need to duplicate such logic for NVAPI support).
* The HLSL source emitter is extended to track whether it emits any operations that require NVAPI support.
* The HLSL source emitter is extended to emit prelude directives based on whether NVAPI is needed and, if it is, to also set the register and space that NVAPI should use based on what was stored in the decoration(s) on the IR module.
* The HLSL source emitter is extended so that it detects global instructions that represent "magic" NVAPI constructs , and emit them as conditional definitions so that they are skipped when NVAPI is included via the prelude.
* The handling of requires capabilities during emit logic was cleaned up a bit so that more logic is shared across targets, and also so that the same logic is used both when emitting a function declaration/definition and when emitting a call to an instrinsic function (which won't get declared/defined).
|
|
* Front-load cuda module loading to fill in RTTI pointers.
* Enable dynamic dispatch codegen for CUDA.
|
|
A long-standing problem for the Slang implementation has been that some targets (notably GLSL/SPIR-V) do not support treating resources (textures, buffers, samplers, etc.) as first-class types. Resource types on such platforms are restricted so that they may not be used as the type of:
1. fields of aggregate types (`struct`s)
2. local variables
3. function results or `out`/`inout` parameters
Issue (1) is handled by our "type legalization" pass today, by splitting aggregates that contain resources into separate fields/variables/parameters. Issue (2) is worked around by putting code into SSA form and promoting local variables to SSA temporaries when possible; the net result is that many local variables of texture type are eliminated (that pass is not perfect, though, and it is possible for users to get errors when it doesn't fully clean up local variables of texture type).
Issue (3) is a much more complicated matter, and it is what this change is concerned with.
A typical solution to issue (3) is to simply inline all of the code in a program, at which point function results and `out`/`inout` parameters will no longer exist to cause problems. We reject such solutions for two reasons. First, there are limitations on control-flow structure in HLSL/GLSL/SPIR-V that mean they cannot express certain programs after inlining has been performed. Second, and more importantly, the philosophy of the Slang compiler is to perform as little duplication of code as possible, so that we do not accidentally contribute to binary size bloat.
Instead, this change tackles the problem of functions that output resource types by adding a new specialization pass. The pass detects functions that ought to be specialized (because they have resource-type outputs), and inspects their bodies to see if the values they output have a predicatable structure that can be replicated outside of the function body. The same logic that inspects the function body also rewrites (a copy of) the function to not have the offending outputs. Finally, all the call sites to a function that is rewritten in this way also get rewritten so that instead of using output values from the function itself, they reproduce the expected output value(s) in their own code.
The pass as presented here is intentionally limited in the scope of what it can optimize away (and the test case only touches on that specific functionality). The goal is to get a basic version of this pass in place and evaluated, and then to expand on its functionality incrementally over time.
|
|
* Allow mixing unspecialized and specialized existential parameters.
* Fixes.
|
|
* Enable lower-generics pass universally.
* Exclude builtin interfaces and functions from lower-generics pass.
* Update stdlib.
* Fixup.
* Fixes handling of nested intrinsic generic functions.
* Fixes.
* Fixes.
|
|
* Tuple types.
* Fix x86 warning
* Improved deduplication
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* AnyValue based dynamic code gen
* Fix aarch64 build error
|
|
The Big Picture
===============
Given input Slang code like:
```hlsl
Texture2D gA;
[shader("compute")]
void kernelFunc(uniform Texture2D b, uint3 tid : SV_DispatchThreadID)
{ ... }
```
the existing CUDA code generation strategy would always generate a kernel with a signature like:
```c++
struct GlobalParams { Texture2D gA; }
struct EntryPointParams { Texture2D b; }
extern "C" __global__
void kernelFunc(EntryPointParams* entryPointParams, GlobalParams* globalParams)
{ ... }
```
This choice was consistent with the conventions of the CPU kernel target, and shares the advantage that it is easy for the user to data-drive the logic for filling in parameters and then invoking a kernel.
However, the approach outlined above has two serious problems when used for CUDA kernels:
* First, it defies the programmer's expectation about what an "equivalent" CUDA kernel signature would be, which makes it awkward for a developer to invoke this kernel from CUDA C++ host code (especially in the context of an app that might also run hand-written CUDA kernels).
* Second, the performance of this approach suffers because every access to a global or entry point parameter turns into a load from global memory. In contrast, a typical hand-written CUDA kernel passes its parameters via an implementation-specific path that (for current CUDA platforms) seems to be equivalent to `__constant__` memory in performance.
This change alters the convention so that the Slang compiler takes the code from the top of this message and translates it into something like:
```c++
struct GlobalParams { Texture2D gA; }
__constant__ GlobalParams SLANG_globalParams;
extern "C" __global__
void kernelFunc( Texture2D b )
{ ... }
```
This translation alleviates both problems with the current translation:
* The signature of the generated CUDA kernel function is as close to that of the original as is possible (we had to eliminate the `SV_*`-semantic varying inputs), and should directly match what the programmer would expect in common cases.
* Entry-point parameters are passed via CUDA kernel parameters, and should thus match in performance. Global parameters are passed via a variable in `__constant__` memory, and thus should also perform as well as possible/expected.
Detailed Changes
================
* Disable the `collectEntryPointUniformParams` pass for CUDA, so that entry-point `uniform` parameters are *not* bundles into a single `struct` and/or `ConstantBuffer`.
* When targeting CUDA, disable the logic for generating an entry-point parameter for passing in the global shader parameter(s)
* Allow `CLikeSourceEmitter` subclasses to override the name generated for entry-point symbols, and use this to add the required prefix for each OptiX kernel type when translating a ray-tracing kernel.
* Add logic to emit "parameter groups" in a specialized way for CUDA (this is the same approach that allows us to generate `cbufffer { ... }` declarations for fxc). A global-scope parameter group will turn into a global `__constant__` variable called `SLANG_globalParams` (that name becomes part of the ABI for Slang-compiled shaders).
* Update the logic in `render-test` for loading and invoking CUDA kernels to handle the new policy.
The last bullet there merits expansion, since it is indicative of the work a client using Slang would have to go through to use our generated kernels with the new policy:
* When loading a CUDA module with one or more kernels, we also use `cuModuleGetGlobal` to query the address of the `SLANG_globalParams` symbol in that CUDA module. That pointer needs to be used when setting global parameter values to be used by kernels in that CUDA odule.
* Because our existing `BindPoint` logic for CUDA always sets up parameter data in GPU memory, we end up having to copy the entry-point parameter data from GPU memory to host memory. This step would ideally be skipped in a codebase that understands the correct policy, but it is a bit unfortunate that it is no longer trivially correct for an application to store all parameter data in GPU memory.
* Before invoking the kernel, we need to use a `cudaMemcpyAsync` to copy from the prepared GPU memory for global parameters over to the `SLANG_globalParams` symbol associated with the kernel to be invoked. Because this operations is issued on the same CUDA stream as the kernel call, it is guaranteed to not overlap with GPU kernel execution.
* When invoking the kernel, we take advantage of the seldom-used `CU_LAUNCH_PARAM_BUFFER_POINTER` facility to specify a contiguous memory region with all the entry-point parameters in it instead of passing each entry-point parameter separately. Given Slang reflection it is also possible to query the offset of each entry-point parameter in the buffer, so we could invoke the kernel in the traditional fashion as well. The choice here is up to the application.
Caveats
=======
* This is a breaking change, and any subsequent release will need to reflect that fact. Any customers who rely on Slang's current CUDA codegen strategy are likely to be surprised by this change, and I don't see an easy way to give them a more gentle transition.
* This change does *not* remove the logic that introduces a `KernelContext` type for code that requires it. That means that things like `static` global variables can continue to work on CUDA for now, but we know that those are not going to be something we can support in the long-term with separate compilation.
* While the policy implemented in this change is a reasonable default, it is still not going to perfectly match expecations for some developers. In particular, some developers who are familiar with both D3D and CUDA will likely wonder why a global `cbuffer` in Slang translates to a global-memory pointer in the output CUDA instead of one global `__constant__` variable per `cbuffer`. A more detailed alternate translation would generate a distinct global `__constant__` variable for each top-level constant buffer or parameter block. We may need to refine the translation even more based on feedback from users who care about how we handle global-scope parameters.
* Recent changes in Slang have broken the logic that handles the OptiX "shader record" as an alternative mechanism for passing entry-point parameters. In order to get any level of OptiX support up and running we will have to change the IR passes that run on CUDA kernels to actually run the "collection" of `uniform` parameters for ray tracing stages, and then to replace references to the resulting parameter with a call to the function to access the shader record.
* The use of `SLANG_globalParams` here works well enough in the case of whole-program compilation; every `CUmodule` ends up with (zero or) one parameter with this name, and an application can just hard-code it. As a mechanism it wouldn't work in the presence of separately-compiled modules that might introduce their own global parameters (including cases like constant lookup tables that really want to be at the global scope). An alternative approach would have Slang generate output PTX for each module, where a module has an optional global symbol for its own global-scope parameters (with a mangled name that is based on the module name), and then a linked CUDA binary has all of those distinct symbols. Such an approach would be compatible with module-at-a-time reflection and parameter binding, but would lead to another breaking change down the line for code that switches to `SLANG_globalParams`.
|
|
* Ensure labels are dumped in `lower-to-ir`.
There is a `dumpIR` function that accepts a label parameter already in slang-emit.cpp. This change moves it to slang-ir.cpp so it may be called from other files.
* update expected test result
Co-authored-by: Yong He <yhe@nvidia.com>
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Run array specialization in a sperate pass.
* rename specializeFunctionCall->specializeFunctionCalls
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Multiple Entry Point Backend
This PR introduces changes to the IR linking, emitting, and options for
multiple entry points. Specifically, this PR updates several locations
to support a (potentially empty) list of entry points, adding list infrastructure and looping over entry points as appropriate.
* Formatting change
* Updated unknown target case to not require an entry point
* Formatting and list consts updates
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Remove KernelContext wrapper from CPU/CUDA emit
Currently, the CPU and CUDA C++ targets rely on a `KernelContext` type that is generated during emit, as a way to provide implicit access to things that were global in the input Slang code, but that can't actually be emitted as globals in the target language (because the semantics of global declarations differ).
For example, input like:
```hlsl
ConstantBuffer<Stuff> gStuff; // shader parameter
groupshared int gData[1024]; // thread-group shared variable
static int gCounter = 0; // "thread-local" global-scope variable
void subroutine() { ... }
[shader("compute")] void computeMain() { ... }
```
would translate to output C++ for CPU a bit like:
```c++
struct KernelContext
{
ConstantBuffer<Stuff> gStuff;
int gData[1024];
int gCounter = 0;
void subroutine() { ... }
void computeMain() { ... }
};
```
Note that both `computeMain()` and `subroutine()` are non-`static` members functions on `KernelContext`, so they have an implicit `this` parameter of type `KernelContext`, which allows the bodies of those functions to implicitly reference `gStuff`, etc. by name in their bodies.
Because `KernelContext::computeMain()` is a member function, we end up emitting an additional global-scope function to expose the entry point to the outside world, and that function is responsible for declaring a local `KernelContext` and invoking the generated entry point on it.
This approach has several important drawbacks:
* It complicates the emit logic for CPU and CUDA, with many special cases around when/how things get emitted
* It complicates the implementation of dynamic dispatch, because what seems like a function pointer in Slang IR needs to be a pointer-to-member-function in C++.
* It makes it difficult to have a non-kernel-oriented mode of compilation for CPU where a Slang function with a given signature gets output as a C++ CPU function with the "same" signature (not wrapped up as a member function of `KernelContext`.
This change makes a step toward addressing these issues by making the introducing of the `KernelContext` type be something that is done in an explicit IR pass instead of being handled as part of the last-mile emit logic.
The most important change is the removal of code related to `KernelContext` from the `slang-emit-{cpp,cuda}.{h,cpp}` files, with the equivalent logic instead being handled in a new pass in `slang-ir-explicit-global-context.{h,cpp}`. It should be noted that further cleanups to the emit logic should now be possible; in particular, both the CPU and CUDA emit paths are manually sequencing the `EmitAction`s instead of relying on the default logic, but at this point they should be able to just use the default. The additional cleanups are left for future work.
The explicit IR pass does more or less what one would expect: it identifies global-scope entities (global variables and parameters) that need to be wrapped and turns them into fields of a `KernelContext` type. It then modifies all entry points to initialize a `KernelContext` as part of their startup. Finally, any code that used to refer to the global entities is changed to refer to a field of the context, with the context passed via new function parameters (the new parameter is only added to functions that need it for now).
Transforming global variables into fields of a `KernelContext` type in the IR pass ends up dropping their initial-value expressions (since those were attached as basic blocks on the `IRGlobalVar`). To avoid breaking code that relies on global-scope (but thread-local) variables, this change also adds an explicit pass that takes the initialization logic on all global variables and moves it to explicit logic that runs at the start of every entry point in a linked module (`slang-ir-explicit-global-init.{h,cpp}`). This pass would also be useful when we get back to direct SPIR-V emit, since SPIR-V also requires initialization logic for globals to be emitted into entry points.
One complication that arises when the IR is introducing the types for entry-point parameters, global-scope parameters, and the `KernelContext` type is that it becomes harder for the emit logic to utter the names of those types (they might not even have names, since `IRNameHint`s might get stripped). This created a problem since the wrapper operations that were being generated for CPU were taking `void*` parameters and casting them to the appropriate type. To work around this issue, we have added an explicit IR pass (`slang-ir-entry-point-raw-ptr-params.{h,cpp}`) that transforms the signature of entry points so that any pointer parameters instead become raw pointer (`void*`) parameters, with the casting being handled inside the entry point itself.
One consequence of all the above changes is that for the CUDA target we no longer need a wrapper function to invoke the generated entry point any more, because the IR function for the entry point ends up having the correct/expected signature already. This is also the case for CPU when it comes to the `*_Thread` wrapper function, but this change doesn't try to eliminate the wrapper because of a belief that the `*_Thread`-level interface is going away anyway.
Because the IR is now responsible for ensuring the signature of the IR entry point for CUDA and CPU is what is expected, I needed to modify the `slang-ir-entry-point-uniforms` pass to always create an explicit parameter for the entry point uniforms when compiling for CUDA/CPU, even if there were no `uniform` parameters on the entry point as written. This also ended up requiring some tweaks to the parameter layout logic to ensure that CPU/CUDA targets always treat `ConstantBuffer<T>` as a `T*` even in the case where `T` is an empty `struct` type (which happens when we construct a `struct` type to represent the uniform parameters of an entry point with no uniform parameters...).
There are several future changes that can/should build on this work:
* We should change the generated signatures for CUDA kernels, so that they don't rely on `KernelContext` for global-scope parameters. At that point we can avoid generating a `KernelContext` at all for CUDA, except when a program uses global-scope thread-local variables.
* We should figure out how to make the "ABI" for dynamic-dispatch calls ensure that the kernel context is either always passed, or always *not* passed. Making a hard-and-fast rule as part of the calling convention for dynamic calls would ensure that they access through the context continues to work with dynamic calls (this change might break it in some cases).
* We should figure out how to handle the layout for the `KernelContext` in cases where a program is composed of multiple separately-compiled modules. Right now the layout of the `KernelContext` requires global knowledge (as does the pass that introduces explicit initialization for global-scope thread-locals).
* We should try to further clean up the CPU/CUDA C++ emit logic to fall back on the default emit behavior more, now that the various special-case approaches that were taken are no longer needed
* fixup: restore build files to default configuration
|
