slang.git/source/slang/slang-ir-legalize-varying-params.cpp, branch master

Allow entry points with missing numthreads on CPU targets (#8678)

2025-10-10T23:01:51+00:00

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.)

Allow 1D SV_DispatchThreadID in CPU targets (#8612)

2025-10-08T23:13:27+00:00

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)

Rename some symbols related to pointers types (#8592)

2025-10-03T04:48:11+00:00

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

[CBP] Pointer frontend changes + groupshared pointer support (#7848)

2025-08-29T22:52:34+00:00

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>

Handle SV_Barycentrics on metal (#8163)

2025-08-13T20:45:44+00:00

Fixes #6785

Remove the semantic decoration from the original entry struct (#8146)

2025-08-13T15:37:22+00:00

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>

Replace [KnownBuiltin] string-based comparisons with enum-based system (#7714)

2025-07-15T10:11:39+00:00

* 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 --------- 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 Co-authored-by: Gangzheng Tong

Legalise out parameters for vertex shaders on metal (#6943)

2025-06-10T09:33:38+00:00

* 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) (#6998)

2025-06-06T20:13:43+00:00

* 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>

Map `SV_VertexID` to `gl_VertexIndex-gl_BaseVertex`, add `SV_Vulkan*ID` semantics (#7150)

2025-05-19T18:21:25+00:00

* 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