| Age | Commit message (Collapse) | Author |
|---|
|
Several tests have compute entry points without a `[numthreads(x,y,z)]`
decoration. Currently, none of these tests run on the CPU target, as
they crash the compiler. I took a look at the SPIR-V emitter, which
falls back to a workgroup size of (1,1,1):
https://github.com/shader-slang/slang/blob/1e0908bd7107dfbdac912b693c3ab9bd6e1dc8b3/source/slang/slang-ir-spirv-legalize.cpp#L1635-L1643
To match this behaviour, this PR implements a fallback solution that
makes `emitCalcGroupExtents()` emit (1,1,1).
This PR is both a question and a suggestion; I'm not sure the approach
here is at all reasonable. Personally, I'd just like to explicitly add
`[numthreads(1,1,1)]` to all such tests, but I don't know if it's
actually legal and supported to not have a `numthreads`. So the
implementation here is a bit conservative.
I ran across these when I went through tests for the upcoming LLVM
target. These were the final blockers to get all autodiff and
language-features tests passing (not counting the ones using things like
wave intrinsics and barriers etc.)
|
|
The varying param legalization pass didn't deal with this 1D form of
SV_DispatchThreadID for CPU targets:
```slang
void computeMain(int i : SV_DispatchThreadID)
```
Instead, it just overrode the type of `i` with a `uint3`, breaking lots
of code that attempted to use `i` for something, like a `switch`
statement for example.
I ran across this when going through `language-feature` tests for the
LLVM target, which will also use this legalization pass. I'm separately
submitting this now because this also fixes the existing CPU target. The
test I enable in this PR is one that was previously generating broken
code on CPU.
(somewhat related issue: #7468)
|
|
Note that while this change touched a large numer of files, there are no
changes to functionality being made here. The only things being done are
renaming various symbols and, in a few cases, updating or adding
comments for consistency with the new names.
The core of the naming changes are:
* Most things named to refer to `OutType` (e.g., `IROutType`,
`IRBuilder::getOutType()`, etc.) have been consistently renamed to refer
to `OutParamType`, to emphasize that the relevant AST/IR node types are
only intended for use to represent `out` parameters.
* The same change as described above for `OutType` is also made for
`RefType`, which becomes `RefParamType` in most cases. One mess that
this exposes is the way that the `ExplicitRef<T>` type in the core
module currently lowers to `IRRefParamType`. This change sticks to the
rule of not making functional changes, so that mess is left as-is for
now.
* Names referring to `InOutType` have been changed to instead refer to
`BorrowInOutType`. The intention with this naming change is to emphasize
that the Slang rules for `inout` are semantically those of a borrow (or
at least our interpretation of what a borrow means).
* Names referring to `ConstRefType` have been changed to instead refer
to `BorrowInType`. This change starts work on clarifying that the
existing `__constref` modifier was never intended to be a read-only
analogue of `__ref`, and instead is the input-only analogue of `inout`.
* The `ParameterDirection` enum type has been changed to
`ParamPassingMode`, to reflect the fact that the concept of "direction"
fails to capture what is actually being encoded, particularly once we
have modes beyond simple `in`/`out`/`inout`.
While this change does not alter behavior in any case (the user-exposed
Slang language is unchanged), it is intended to set up subsequence
changes that will work to make the handling of these types in the
compiler more nuanced and correct. Breaking this part of the change out
separately is primarily motivated by a desire to minimize the effort for
reviewers.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Resolves #7628
Resolves: #8197
Primary Goals:
1. Add `Access` to pointer
2. AddressSpace::GroupShared support for pointers (SPIR-V)
3. Add `__getAddress()` to replace `&`
* `&` is not updated to `require(cpu)` since slangpy uses `&`. This
means we must: (1) merge PR; (2) replace `&` with `__getAddress()`; (3)
add `require(cpu)` to `&`
Changes:
* Added to `Ptr` the `Access` generic argument & logic (for
`Access::Read`).
* Moved the generic argument `AddressSpace` from `Ptr` to the end of the
type.
* Added pointer casting support between any `Ptr` as long as the
`AddressSpace` is the same
* Disallow globallycoherent T* and coherent T*
* Disallow const T*, T const*, and const T*
* Fixed .natvis display of `ConstantValue` `ValOperandNode`
* Support generic resolution of type-casted integers
* Added `VariablePointer` emitting for spirv + other minor logic needed
for groupshared pointers
Breaking Changes:
* Anyone using the `AddressSpace` of `Ptr` will now have to account for
the `Access` argument
* we disallow various syntax paired with `Ptr` and `T*`
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Fixes #6785
|
|
When we legalize the entry point param, there are cases where we need to
reconstruct a struct for the parameter and the original struct wouldn't
be used. But if the user tries to use the origianl struct as a type for
a function parameter, we will end up using both the original struct and
the synthesized struct at the same time.
On Metal and WGSL, it causes an error when an identical semtaic is used
on more than one variable.
This commit removes the semantics from the original struct after cloning
the type.
Fixes https://github.com/shader-slang/slang/issues/8141
Related to https://github.com/shader-slang/slang/issues/7693
---------
Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
|
|
* Initial plan
* Implement enum-based KnownBuiltin system to replace string comparisons
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Add test for enum-based KnownBuiltin system and verify functionality
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Implement enum-based KnownBuiltin system with direct integer values
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Fix IntVal access and update tests for new enum-based KnownBuiltin system
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Replace hardcoded KnownBuiltin integers with preprocessor enum syntax
- Updated all KnownBuiltin attributes to use $( (int)KnownBuiltinDeclName::EnumValue) syntax
- Added space between parentheses to avoid preprocessor bug: $( (int) instead of $((int)
- Updated both core.meta.slang and hlsl.meta.slang files
- Eliminates preprocessor-time integer conversion, baking enum values directly into meta files
- Maintains same functionality while using type-safe enum references
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Fix IDifferentiablePtr KnownBuiltin mapping regression
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Remove unused IDifferentiablePtrType enum case from KnownBuiltinDeclName
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Clean up temporary AST dump files from testing
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Replace hardcoded integer with descriptive constant in KnownBuiltin test
Replace the hardcoded [KnownBuiltin(0)] with a descriptive named constant
GEOMETRY_STREAM_APPEND_BUILTIN to improve code readability and maintainability.
The test now clearly indicates which builtin enum value is being tested.
🤖 Generated with [Claude Code](https://claude.ai/code)
Co-authored-by: Gangzheng Tong <gtong-nv@users.noreply.github.com>
---------
Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: Gangzheng Tong <gtong-nv@users.noreply.github.com>
Co-authored-by: Gangzheng Tong <tonggangzheng@gmail.com>
|
|
* Handle pointer types when getting type cast style
Closes https://github.com/shader-slang/slang/issues/6025
* Move vertex shader out parameters to return type for Metal
Closes https://github.com/shader-slang/slang/issues/6025
* More asserts
* Make struct instead of tuple
* More layout preservation
* Handle same function result
* more layout
* remove layout
* a
* more debug code
* more debug code
* a
* layout working
* refactored
* more tests
* more tests
* fuse loops
* remove unused comments
* Correct filecheck usage
* debug code
* correct name and order of filecheck vars
* simplify
* Address review comments
fix warning
* simplify handling of simple vertex shaders
|
|
* Address issues with GLSL style global in/out vars (#6669)
Asserts and segfaults were observed trying to compile a simple
vertex shader like:
````
in int2 inPos;
[shader("vertex")]
main(uniform int2 test1, int2 test2, out float4 pos: SV_Position)
void main()
{
// Bogus use of all input vars to prevent optimizing out.
pos = float4(inPos.x, test1.x, test2.y, 0);
}
````
Further investigation found that while replacing "uniform int2 test1"
with "int2 test1" allowed for successful compilation, the resulting
output shader would have overlapping location qualifiers. For example,
compiling the above with "int2 test1" to glsl might give:
````
...
layout(location = 0) in ivec2 test1_0;
layout(location = 1) in ivec2 test2_0;
layout(location = 0) in ivec2 translatedGlobalParams_inPos_0;
...
````
This was because Slang does not actually support mixing GLSL style global
in/out vars and entry point params. However, this is never checked for
or noted in documentation. Slang source also assumes input shaders do not
mix these and these assumptions ultimately led to the observed asserts
and seg faults when using uniform entry point params.
This change makes updates to throw an error when the compiler detects that
it is trying to translate global in/out variables into entry point params
when an entry point already contains parameters, allowing for compilation
to fail gracefully.
Certain tests have been updated to avoid mixing GLSL style global in/out
vars and entry point params. This was mostly for tests that were using
functions like WaveGetLaneIndex which use global in vars for certain
platforms (see __builtinWaveLaneIndex).
* Address issues with GLSL style global in/out vars - updates 1 (#6669)
Update addresses review feedback to support mixing GLSL-flavored global
in/out vars and entrypoint parameters when either all global in/out vars
or all entry point params have a system value binding semantic.
* Address issues with GLSL style global in/out vars - updates 2 (#6669)
This update attempts to actually allow mixing GLSL style global in
vars and entry point vars.
Change attempts to recalculate offsets when adding the global input
vars into the recreated entry point params layout.
Additional updates were made to:
-resolve further issues uncovered with entry point uniform params.
-Address improper use of SV_DispatchThreadID in wave-get-lane-index.slang
for metal. "thread_position_in_grid" is not supported for signed integer
scalars or vectors.
-Fix a spirv casting conflict due to the implementation of
gl_PrimitiveID.get conflicting with PrimitiveIndex().
-Add a call to remove a global var in replaceUsesOfGlobalVar(). The global
var is already replaced in this function and keeping it around can prevent
it from being cleaned up by DCE if it still has decorations.
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
semantics (#7150)
* Map SV_VertexID to `gl_VertexIndex - gl_BaseVertex`, provide SV_Vulkan* SV semantics
* Fix docs
* Regenerate toc
* Fix affected pointer-2 test
* Add tests
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* Add noreturn attribute to IgnoreHit
* Revert "Add noreturn attribute to IgnoreHit"
This reverts commit 3cf2354dada71b9a8713b08f3a2e261de4aabfa4.
* Fix: Preserve inout param modifications with OptiX IgnoreHit()
Issue #6326 identified that in OptiX, when using IgnoreHit() (which
maps to the "noreturn" optixIgnoreIntersection() intrinsic), any
modifications made to 'inout' parameters within the shader would be
lost. This was due to IgnoreHit() preventing the execution of the
copy-back operation from the temporary variable (used to implement
'inout' semantics) to the original parameter.
This commit introduces a new IR pass, 'undoParameterCopy', specifically
for CUDA/OptiX targets to address this. The pass operates as follows:
1. Identifies temporary IR variables created for 'inout' parameters,
which are now decorated with 'TempCallArgVarDecoration'.
2. Maps these temporary variables back to their original parameter
storage (e.g., the OptiX payload pointer).
3. Replaces all uses of the temporary variable directly with the
original parameter pointer.
4. Removes the temporary variable declaration and its initializing store
(which copied from the original parameter to the temporary).
By transforming the IR to operate directly on the original parameter
storage before any potential call to IgnoreHit(), this fix ensures
that all modifications are preserved, correctly resolving issue #6326.
The pass is integrated into the compilation flow for relevant targets.
* Refactor(IR): Optimize GetOptiXRayPayloadPtr for better DCE/CSE
To allow for more effective dead code elimination (DCE) and
common subexpression elimination (CSE) of `getOptiXRayPayloadPtr`
instructions, this commit:
- Marks `kIROp_GetOptiXRayPayloadPtr` as side-effect-free within
`IRInst::mightHaveSideEffects` (in `slang-ir.cpp`).
- Flags `GetOptiXRayPayloadPtr` as `HOISTABLE` in its definition
within `slang-ir-inst-defs.h`.
This addresses scenarios where multiple, potentially redundant,
calls to `getOptiXRayPayloadPtr` might appear in the IR,
allowing optimizers to produce cleaner and potentially more
efficient code for OptiX targets. This change supports efforts
to refine IR handling for ray-tracing shader stages.
* Remove debugging code
* Refactor UndoParameterCopyVisitor for improved performance
- Optimized IR traversal by combining multiple passes into a single scan
- Removed unnecessary dictionary, immediately replace uses when a temp var is found
- Reduced duplicate code paths by checking for both temp vars and redundant stores in one loop
- Better handling of the 'changed' flag to ensure DCE only runs when needed
- Results in fewer instruction traversals and improved efficiency for large functions
* Add Test
|
|
* cuda: Improve entry handling for SV_DispatchThreadID
Fixes #6780
This commit improves CUDA entry point handling by extracting appropriate
components from DispatchThreadID based on parameter type. It now
properly handles uint scalar (x component only) and uint2 vector (x,y
components) instead of always using the full uint3 value.
Add a new test case to check for this.
* format code
* fix CI failure
* Handle review comments
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
|
* Add gl_PointCoord support in GLSL compat mode
* Add SV_PointCoord
* Test on metal as well
* Update SPIRV system value semantics table in docs
* Update metal docs for SV_PointCoord
|
|
|
|
* initial wip
* more wip
* add test
* add unexpected for invalid target
* fixups and improve error message
* fixups and improve error message
* remove incorrect comment
---------
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
|
* support WaveGetLane* for WGSL and Metal
* update test and glsl support
* address review comments and fix metal test
* add missing pragma guard
* update test
* Revert "update test"
This reverts commit f2b97e91c29de154190710580c343bd0764aedbb.
* update failing glsl metal test and added new test
* make hlsl and glsl outputs similar
* update test
* disable tests for Metal and cleanup
* comment fix
* add expected failures
* correct expected failures list
* remove expected failure
* add tests to expected failure
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* Refactor to reuse common for metal and wgsl entry point legalization
* refactor system val work item
* refactor simplify user names
* clean up fix semantic field of struct
* improve code layout
* split wgsl/metal to seperate classes and cleanup
* remove extra includes
* remove dead code comments
* minor cleanup
* squash merge from master and resolve conflict
* apply metal spec const thread count changes
* Revert "apply metal spec const thread count changes"
This reverts commit c42d707fd25ee0328598650d3235cd2322810ccc.
* Revert "squash merge from master and resolve conflict"
This reverts commit 06db88ef7001bdfe93fb23af35af0d026b255dee.
* Merge remote-tracking branch 'origin/master'
* apply metal spec const thread count changes
* Revert "apply metal spec const thread count changes"
This reverts commit 3b9e6f53cee2e6076ac2b7a0d015a1ed2cbbd627.
* Revert "Merge remote-tracking branch 'origin/master'"
This reverts commit 99869d573a46dadeb24445405f5a1e37a8e03d0d.
* apply metal spec const thread count changes
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* Fix optix varying legalization.
* Add test.
|
|
* Allow using specialization constants in numthreads attribute
* Add support for GLSL local_size_x_id syntax
* Fix overeager specialization constant parsing
* Add diagnostics for specialization constant numthreads
* Remove unused variable
* Fix local_size_x_id not finding existing specialization constant
* Allow materializeGetWorkGroupSize to reference specialization constants
* Use SpvOpExecutionModeId for modes that require it
* Cleanup specialization constant numthreads code
* Add tests for specialization constant work group sizes
* Fix implicit Slang::Int -> int32_t cast
* Fix querying thread group size in reflection API
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* Add executable test on matrix-typed vertex input.
* Fix emit logic of matrix layout qualifier.
* Pass fragment shader varying input by constref to allow EvaluateAttributeAtCentroid etc. to be implemented correctly.
|
|
* Move switch statement bodies to their own lines
* format
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* format
* Minor test fixes
* enable checking cpp format in ci
|
|
* initial change to test with CI for CPU/CUDA errors
* Fixes to Metal Input parameters and Output values
Note:
1. Flattening a struct is the process of making a struct have 0 struct/class members.
Changes:
1. Separated `legalizeSystemValueParameters`. This was done to make it easier to run `legalizeSystemValue` 1 system-value at a time to simplify logic. This change is optional and can be undone if not preferred.
2. Wrap everything inside a Metal legalization context. This was done since it simplifies a lot of logic and will be required for #4375
3. Created `convertSystemValueSemanticNameToEnum` and expanded the existing System-Value Enum system. This allows (sometimes) faster comparisons and helps prepare code for porting into `slang-ir-legalize-varying-params.cpp` (#4375)
4. Added a more dynamic `legalizeSystemValue` system so more than 2 types can be targeted for legalization. This is required to legalize `output`. There is still no preference for any converted type, the first valid type will be converted to.
5. Flatten all input(`flattenInputParameters`)/output(part of `wrapReturnValueInStruct`) structs and assign semantics accordingly.
6. Semantics when legalized have no specific logic other than to: 1. avoid overlapping semantics 2. Prefer assigning explicit semantics specified by a user.
7. Fixed some issue with incorrect output semantics if not a fragment stage (when there are not any assigned semantics)
* change metallib test to the correct metal test
* comment code & cleanup -- Did not address all review
Added comments for clarity + cleaned up some odd areas which were messy
* Add comment to `fixFieldSemanticsOfFlatStruct`
I found `fixFieldSemanticsOfFlatStruct` to still be confusing at a cursory glance. Added comments to make the function be more understandable.
* white space
* Address review comments
1. Fix semantic propegation.
2. Fix how we map struct fields of the flat struct to struct. This is specifically important for if reusing the same struct twice since struct member info is not unique per struct instance used.
* Fix semantic legalization by adding TreeMap
Add TreeMap to allow proper sorted-object data iteration.
* Fix some compile issues
* try to fix gcc compile error
* compile error
* fix logic bug in treeMap iterator next-semantic setter
* fix vsproject filters
* filter file syntax error
* remove need of a context to make copies stable
* Rename treemap to the more appropriate name of "treeset", adjust code comments accordingly.
* remove custom type `TreeSet` and use `std::set`
* remove TreeMap fully
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
* Implement for metal `SV_GroupIndex`
1. If we don't have `sv_GroupThreadId` available we create one using `SV_GroupIndex`s location data.
2. We emit code emulating `sv_GroupThreadId` from the same logic that CUDA/CPP uses.
* address most review comments
Addressed all but two: [1](https://github.com/shader-slang/slang/pull/4385#discussion_r1639058473) and [2](https://github.com/shader-slang/slang/pull/4385#issuecomment-2166934855)
I want to enable tests and be sure there is no bugs using CI before I redesign the code so I have a working fallback.
* address comment, enable tests
enable now functioning tests due to `SV_GroupIndex` working with metal
* syntax error with groupThreadID search
did `= param` instead of `= i.param`
* add `sv_groupid` for test + test fixes
* disable test that won't work regardless
|
|
|
|
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* Clean up `IRReturnVoid`.
* Update gitignore.
Co-authored-by: Yong He <yhe@nvidia.com>
|
|
* Cleanup refactoring work around the IR builder
We have some long-term goals for the IR that require a more centralized and disciplined set of rules for how IR instructions get created/emitted. I had been working on trying to set things up so that all IR instruction creation goes through a single bottleneck point, but the non-trivial work in that branch was getting drowned out by the sheer volume of cleanup and refactoring changes. This change tries to pull together several of the more important cleanups.
The big pieces are:
* `IRBuilder` and `SharedIRBuilder` now protect their data members and rely on users to initialize them more directly via constructor of an `init()` method. This change affects a *bunch* of sites where `IRBuilder`s were created. I changed use sites to use the constructors whenever possible, and to use `init()` in cases where we had longer-lived builders that needed to be initialized multiple times.
* The insertion location for the `IRBuilder` now uses an encapsulated type called `IRInsertLoc`. This new type can replace what used to be just two `IRInst*` fields in the builder, and also covers some new functionality (if we ever want to take advantage of it). Very little client code cares about this change, but it is still a nice cleanup in terms of making things more explicit.
* The creation of an `IRModule` has been moded *out* of `IRBuilder`, because in practice we `IRBuilder` always wants to be associated with a pre-existing `IRModule` at creation time (via its `SharedIRBuilder`). There is now an `IRModule::create()` operation instead. This required changing the sequencing at many `IRModule` creation sites, since most had been contriving to make an `IRBuilder` first. There were also several cleanups because code had been carelessly using non-reference-counted pointers for `IRModule`s in ways that broke now that `IRModule::create()` always returns a `RefPtr`.
* The core operations to actually allocate memory for IR instructions were moved into `IRModule` (since they interact with the memory pool that the module owns). These *were* called `createEmptyInst()` but have been renamed into `_allocateInst()`. In principle these seem like they should only be needed to be called by the `IRBuilder`, but in practice they are also needed by the IR deserialization logic.
* A few core operations for emitting IR instructions that were associted with `IRBuilder` were moved to actually be methods on `IRBuilder`. First is `_findOrEmitConstant` which is the primary bottleneck for creating simple scalar constant values. Another is `_createInst` (formerly part of the templated `createInstImpl` along with `createInstWithSizeImpl`) which is the main bottleneck for allocation and initialization of any instruction other than a constant (well, the `IRModuleInst` is the other exception...). Finally, there is also `_maybeSetSourceLoc()`, which is obvious to scope inside the `IRBuilder` once it is protecting the source-location info.
Notes:
* The `minSizeInBytes` parameter to `_createInst()` might not actually be needed at all. At this point any `IRInst` subtypes that need data allocated for things other than their operands already get created manually via `_allocateInst` or `_findOrEmitConstant`, so I *think* we could remove that part. I will handle that in a subsequent cleanup if it turns out to be the case.
* There is one IR pass (`slang-ir-string-hash.cpp`) that is using manual `_allocateInst()` instead of going through an `IRBuilder`. It could be easily cleaned up to not do so (and I will probably make that change down the line), but for now I wanted to avoid doing anything that wasn't close to pure refactoring if I could.
* At this point in our design an `IRBuilder` is a very lightweight thing - it basically just owns the insertion location plus a source location to write into instructions. A lot of our code currently treats `IRBuilder`s like they are expensive and/or need to be re-used (which leads to them being used in more mutable/stateful ways). It is quite likely that as we clean up other aspects of the implementation of IR creation/emission we can make `IRBuilder` use feel more lightweight in ways that can streamline and simplify code.
* The next step for this work is to identify the different paths that eventually lead to `_createInst()` being called, and unify them at a single bottleneck operation that can own the decisions around when to create an instruction vs. when to re-use an existing one (rather than those decisions being baked into the various `IRBuilder` subroutines that create instructions of the various subtypes).
* fixup: gcc/clang C++ spec details
|
|
* OptiX ray payload can now be read from and written to using the two payload register pointer method
* changing op to more descriptive name
* small tweak to allow for dumping out intermediate source for cuda targets
* initial trace ray call compiling
* hit attributes now work for float and int types, and vectors thereof
* Hitgroups using structs and arrays now work with optix
Co-authored-by: T. Foley <tfoleyNV@users.noreply.github.com>
|
|
* OptiX ray payload can now be read from and written to using the two payload register pointer method
* changing op to more descriptive name
* fixup: comment change to re-trigger CI
Co-authored-by: T. 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
|
|
* 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
|
|
The main change here is that the CPU and CUDA C++ emit paths now rely on an earlier IR pass to legalize the varying parameter list of a kernel and translate references to varying parameters with semantics like `SV_DispatchThreadID`. Doing so removes a lot of special-case logic from the emit passes.
This work moves us even closer to being able to eliminate `KernelContext` from the CPU/CUDA emit logic, because it removes the issue of state related to varying inputs being stored in `KernelContext`.
The new pass that handles the legalization is in `slang-ir-legalize-varying-params.cpp`, and it borrows heavily from the existing `slang-ir-glsl-legalize.cpp` pass. The new pass factors out the target-independent and target-dependent logic, so that both CPU and CUDA can share much of the same code despite having very different rules for how the system-value parameters are being provided.
An eventual goal is to have the new pass also handle the GLSL case, but doing so requires copying even more logic out of the GLSL-specific pass, and doing so seemed like a step to far for what was meant to be a stepping-stone change as part of other work. As a result of the incomplete nature of the pass, certain cases don't work for compute shader inputs for CPU/CUDA (e.g., wrapping your varying inputs in a `struct` type parameter), but those were cases that also didn't work in the existing `emit`-based logic.
One major consequence of this change is that the logic for emitting the various different functions that represent an entry point for our CPU back-end has been streamlined and simplified. The original logic had a fair bit of cleverness built in to try and avoid unnecessary math ops when computing the various IDs/indices, while the new logic is much more simplistic (the main dispatch function loops over threadgroups with a triply-nested `for` and then delegates to the group-level function loops over threads with its own nested `for`s).
Longer term, it will be important to simplify the CPU functions we emit further, by eliminating things like the `_Thread` function that should never really be exposed to users (the minimum granularity of invoking a CPU compute kernel should be a single threadgroup). We may eventually decide to synthesize all of the extra code that is being generated in the `emit` pass as IR instead.
|
