summaryrefslogtreecommitdiffstats
Commit message (Collapse)AuthorAge
* Optimize CapabilitySet deserialization performance (#8552)Ellie Hermaszewska2025-09-29
| | | | | Closes https://github.com/shader-slang/slang/issues/8477 About a 50% reduction in deser performance for capability sets
* cmake: install config files into the correct directory for find_package() ↵David Aguilar2025-09-29
| | | | | | | | (#8482) cmake's find_package() will only find the configs when they are in a slang/ subdirectory. Co-authored-by: Mukund Keshava <mkeshava@nvidia.com>
* Update function type after inserting KernelContext parameter (#8551)Julius Ikkala2025-09-29
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Without this, there are functions with missing parameters in their type in the IR after running the `introduceExplicitGlobalContext` pass: ``` [layout(%15)] [export("_SV4test12outputBuffer")] [nameHint("outputBuffer")] let %outputBuffer : _ = key [noSideEffect] [export("_S4test7dostuffp1pi_ff")] [nameHint("dostuff")] func %dostuff : Func(Float, Float) { block %34( [nameHint("f")] param %f : Float, [nameHint("kernelContext")] param %kernelContext : Ptr(%KernelContext, 0 : UInt64, 1 : UInt64)): let %35 : Float = mul(%f, %f) let %36 : Ptr(ConstantBuffer(%GlobalParams, DefaultLayout), 0 : UInt64, 1 : UInt64) = get_field_addr(%kernelContext, %globalParams) let %37 : ConstantBuffer(%GlobalParams, DefaultLayout) = load(%36) let %38 : Ptr(RWStructuredBuffer(Float, DefaultLayout, %20)) = get_field_addr(%37, %outputBuffer) let %39 : RWStructuredBuffer(Float, DefaultLayout, %20) = load(%38) let %40 : Ptr(Float) = rwstructuredBufferGetElementPtr(%39, 1 : Int) let %41 : Float = load(%40) let %42 : Float = mul(%35, %41) return_val(%42) } ``` Not sure why this doesn't seem to negatively affect existing targets, but it sure is an issue for the LLVM target I'm working on. I could've left this fix for that PR, but I want to check now if this causes any issues with the existing targets using the CI. This also happens with the entry point functions, where the function type is not updated after adding `ComputeThreadVaryingInput`. This had no effect in the C++ target because `convertEntryPointPtrParamsToRawPtrs(irModule);` is called right after and fixes it.
* Fix segfault when shader entry points return resource types (#8434)Copilot2025-09-29
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | The Slang compiler was segfaulting when trying to compile shaders that return resource types (like `Texture2D`, `RWTexture2D`, `SamplerState`, etc.) from entry point functions. This occurred because there was missing validation that should reject such invalid return types before they reach IR generation. For example, this code would cause a segfault: ```slang StructuredBuffer<Texture2D<int>> skyLight; [shader("compute")] Texture2D<int> computeMain(uint3 threadID : SV_DispatchThreadID) { return skyLight[threadID.x]; } ``` ## Root Cause The issue was in the entry point validation logic in `validateEntryPoint()`. While there was a TODO comment indicating that return type validation should be performed, it was never implemented. The compiler would accept the invalid shader code and attempt to process it during IR lowering, where resource types as return values are not properly handled, leading to a segmentation fault. ## Solution 1. **Added robust validation**: Modified `validateEntryPoint()` in `slang-check-shader.cpp` to use the existing `SemanticsVisitor::getTypeTags()` functionality to check for invalid return types by detecting `TypeTag::Opaque` and `TypeTag::Unsized` bits. This leverages the existing type analysis infrastructure that comprehensively handles: - Direct resource types (Texture2D, RWTexture2D, SamplerState, etc.) - Structs containing resource-typed fields (through type tag propagation) - Nested structures and complex type hierarchies - Arrays and other composite types 2. **Added diagnostic message**: Uses existing diagnostic `entryPointCannotReturnResourceType` (error 38010) that provides a clear error message explaining why resource types cannot be returned from shader entry points 3. **Updated existing tests**: Modified existing tests to match the updated validation behavior ## Result Instead of a segfault, users now get a clear, actionable error message: ``` error 38010: entry point 'computeMain' cannot return type 'Texture2D<int>' that contains resource types ``` The fix properly handles all resource types including `Texture2D`, `RWTexture2D`, `SamplerState`, and others, while preserving the ability to compile valid shaders that return simple data types. Fixes #6438. <!-- START COPILOT CODING AGENT TIPS --> --- 💡 You can make Copilot smarter by setting up custom instructions, customizing its development environment and configuring Model Context Protocol (MCP) servers. Learn more [Copilot coding agent tips](https://gh.io/copilot-coding-agent-tips) in the docs. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: expipiplus1 <857308+expipiplus1@users.noreply.github.com> Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com> Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
* Add SPV_NV_bindless_texture support (#8534)Lujin Wang2025-09-26
| | | | | | | | | | | | | Treat DescriptorHandle as uint64_t instead of uint2. Implement target-specific SPIR-V emission with the bindless texture support. For OpImageTexelPointer, Image must have a type of OpTypePointer with Type OpTypeImage. Fix the issue by using [constref] in __subscript. Add a test coverage for various texture/sampler handle types. --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Diagnostic on use of unsupported entry point modifiers (#8487)James Helferty (NVIDIA)2025-09-26
| | | | | | | | | | | | Generate a diagnostic warning whenever unsupported modifiers (keywords, attributes) are found on entry point parameters. These have been silently ignored up until now, with the parser accepting them but Slang not actually doing anything with them. Fixes #7151 --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Remove `cooperative-vector` tests from expected CI failures (#8381)aidanfnv2025-09-26
| | | | | | | | | | | | Fixes #7715 Updating the Vulkan SDK on the Windows CI machines to 1.4.321.1 has fixed some illegitimate VVL errors in the `cooperative-vector` tests, and #8541 has fixed some legitimate VVL errors in some of those tests, so now they can be removed from the list of expected test failures. The only expected `cooperative-vector` failures that remain are for `-emit-spriv-via-glsl`, as we do not support `GLSL_NV_cooperative_vector` yet, see #7727.
* Fix VVL errors on coopvec tests (#8541)Jay Kwak2025-09-25
| | | | | | | | | | | | | | | | | | | | | | It appears that the inputType of the coopvec-mat-mul cannot be signed int32. It could be floating types or signed int32. Changing the tests to use uint32 instead of int32. The spec guarantees the following combinations and the rest should be queried at the runtime if it is supported by the HW. https://registry.khronos.org/vulkan/specs/latest/man/html/VkCooperativeVectorPropertiesNV.html#_description inputType | inputInterpretation | matrixInterpretation | biasInterpretation | resultType -- | -- | -- | -- | -- FLOAT16 | FLOAT16 | FLOAT16 | FLOAT16 | FLOAT16 UINT32 | SINT8_PACKED | SINT8 | SINT32 | SINT32 SINT8 | SINT8 | SINT8 | SINT32 | SINT32 FLOAT32 | SINT8 | SINT8 | SINT32 | SINT32 FLOAT16 | FLOAT_E4M3 | FLOAT_E4M3 | FLOAT16 | FLOAT16 FLOAT16 | FLOAT_E5M2 | FLOAT_E5M2 | FLOAT16 | FLOAT16
* Fix CMake error with generator on list in 3.22.1 (#8364)Dario Mylonopoulos2025-09-25
| | | | | | | | | Fixes #8335 --------- Co-authored-by: Mukund Keshava <mkeshava@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com> Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
* Adding 4 WGPU tests to expected-failure-github (#8540)Jay Kwak2025-09-25
| | | | | | | | | | Four WGPU tests print VVL errors. And it is preventing us from upgrading VulkanSDK on CI machines. This commit put them in the expected-failure-github.txt so that we can continue upgrading VulkanSDK. They will be re-enabled when the following issues are resolved: - https://github.com/shader-slang/slang/issues/8145 - https://github.com/shader-slang/slang/issues/8379
* Prepare VulkanSDK release Oct 2025 (#8525)Jay Kwak2025-09-25
| | | | Related to - https://github.com/shader-slang/slang/issues/8519
* Add timeout to CI jobs; add mimalloc to .gitignore (#8538)Gangzheng Tong2025-09-25
|
* Remove unnecessary Load and Store pair (#8433)Jay Kwak2025-09-24
| | | | | | | | | | | | | | | | | | | | This commit removes unnecessary Load and Store pairs in IR. When the IR is like ``` let %1 = var let %2 = load(%ptr) store(%1 %2) ``` This PR will replace all uses of %1 with %ptr. And the load and store instructions will be removed. But I found that there can be cases where %2 might be still used later in other IRs. For these cases, the removal of load instruction relies on DCE. --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
* Adding slang-test option to ignore abort popup message (#8492)Jay Kwak2025-09-23
| | | | | | | | | | | | | | | | | | | | | | | | | | With the recent Windows runtime libraries, a new popup window started appearing when `abort()` is called. This was observed when VVL prints a message as a part of WGPU test. Although it can be helpful when we want to debug it, it breaks the behavior of CI scripts when the tests are expected to continue even when they fail. When the test fail, CI script stops in the middle and wait for a user to click on a button on the dialog window, which cannot happen. As a result, when there is a VVL error message, CI run stops in the middle and the testing stops prematurely. This commit adds a new command-line argument, `-ignore-abort-msg`, that ignores the abort message and it wouldn't show the dialog popup window. From the implementation perspective, there are three places that are related. - slang-test itself should turn off the flag. - render-test should turn off the flag after getting the argument from slang-test - test-server should turn off the flag after getting the argument from slang-test When test-server runs render-test, the arguments are already handled by slang-test, so test-server needs to just pass through the arguments.
* Enhance SPIRV-Headers update instructionsJay Kwak2025-09-23
| | | Added instructions to retrieve the hash value of spirv-headers.
* Legalize type as well in legalizeOperand (#8483)Gangzheng Tong2025-09-23
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This fixes a type mismatch issue. See the generated cuda code ```cuda struct Query_0 { EmptyExample_0 query_0; uint hasNonEmptyAbsorbingBoundary_0; }; struct Query_1 { uint hasNonEmptyAbsorbingBoundary_0; }; struct GlobalParams_0 { Query_0* gQuery_0; RWStructuredBuffer<float3 > gInput_0; RWStructuredBuffer<float> gOutput_0; }; ... Query_1 _S4 = *globalParams_0->gQuery_0; // ==> type mismatch at call site! ``` **Root Cause:** During the empty type legalization pass in Slang's IR processing, struct types were being optimized. e.g., `Query_0` → `Query_1` with empty type removed), but this created an inconsistency: **Function parameters were updated:** When Query_compute_0 function was legalized, its parameter type was correctly updated from `Query_0` to the optimized `Query_1` **Global parameter types were NOT updated:** The `ParameterBlock<Struct>` type in globalParams still referenced the old `Query_0` type The PR adds special handling for type operands in the `legalizeInst` function. This triggers the legalization of the `StructType` from the original `legalizeOperand` call site. The leaglized result will be saved in the type-to-legal-type map and be re-used when the same type requires legalization again (e.g. in the `IRFunc` as parameter) Fixes: https://github.com/shader-slang/slang/issues/7905
* Lookup refactor (#8467)kaizhangNV2025-09-23
| | | | | | | | | | | Close #8201. This PR unify the lowering logic for LookupDeclRef of an interface requirement. We will always lower this AST node to a LookupWitness IR. The key of this IR is the special witnessTableType `ThisTypeWitness`, this witness Table is simply a wrapper for an interface type. Our current specialization pass doesn't handle this kind of LookupWitness IR at all, so we will also add the specialization of this_type IR as well.
* fix a crash when using type equality constaint (#8515)kaizhangNV2025-09-23
| | | | | | | | | | | | Close #8193. When constructing `TransitiveTypeWitness` node, we should check if there is operand that represents two equal times. Currently, we only check whether the operand is `TypeEqualityWitness`, which is not good enough, because a `DeclaredSubtypeWitness` could also be representing two same types, in that case, we should also const fold this kind of witness. Fails to do so, we could finally ends up with a generating a lookup witness IR on a generic parameter that is not supposed to be looked up.
* Fix varying output structs in GLSL source (#8501)Julius Ikkala2025-09-23
| | | | | | | | | | | | | | | | | | | | | | | | Closes #8500. `slang-ir-translate-global-varying-var.cpp` turns the global varying outputs into a struct that's returned from the entry point. Currently, there's a problem when one of the outputs is a struct. It always creates a generic `IRTypeLayout`, even when a correct type layout already exists. Somehow, this appears to work when the global varying outputs aren't structs. The crash occurs in `slang-ir-glsl-legalize.cpp:createGLSLGlobalVaryingsImpl()`. It correctly handles the generated outer struct, but when that contains an inner struct, it's been given a non-struct type layout and crashes. This PR uses the correct layout if found, instead of generating a broken placeholder. This matches the behaviour that has already been implemented for inputs. Additionally, I removed a call to `addResourceUsage` from both the input and output side. I can't see any way in which it would've affected anything, the layout builder is never used after that call and it doesn't retroactively modify the layout that was already created.
* Split overloaded uses of RefType in front-end (#8427)Theresa Foley2025-09-23
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Overview ======== This change is the start of an attempt to address how the Slang compiler codebase has ended up conflating two similar, but semantically distinct, concepts: * The long-standing notion of `ref` parameters (only allowed for use in the builtin modules), which are encoded using a wrapper `Type` in the AST as part of the representation of the parameters of a `FuncType`. * A recently-introduced notion of explicit reference types that mirror the built-in `Ptr` type, with a relationship comparable to that between pointer and reference types in C++. The change splits the `Ref<T>` type in the core module into two distinct types, with one for each of the two use cases. Similarly, the `RefType` class in the compiler's AST is split into two distinct classes, to represent the two cases. Background ========== The `Ref<T>` type in the core module (hidden and not intended for users to ever see or use) was originally introduced to encode the `ref` parameter-passing mode, comparable to the hidden `Out<T>` and `InOut<T>` types used to encode `out` and `inout` parameter-passing modes. The `Ref<T>` type in the core module was encoded as a instance of the `RefType` class in the Slang AST (similar to how `Out<T>` mapped to an `OutType`). These AST classes were *only* intended to be used by the compiler front-end as part of its encoding of function types. The `FuncType` class needed a way to distinguish an `inout int` parameter from a plain (implicitly `in`) `int` parameter, so these wrapper like `RefType` and `OutType` were introduced to encode both the parameter type (`T`) and the parameter-passing mode in a form that could be passed around as a `Type`. Notably, the `Ref<T>` type (and `Out<T>`, etc.) were *not* intended to be type names that ever get uttered in Slang code (not even in the builtin modules), and the vast majority of the compiler code was not supposed to ever encounter them. They were an implementation detail of `FuncType`, and nothing else. (In hindsight it may have been a mistake to use a nominal type declared in the core module to implement these wrappers; it might have been a good idea to use an entirely separate class of `Type` for this case...) Recent changes to the builtin modules introduced functions that wanted to *return* a reference (so that the parameter-passing-mode modifiers like `ref` could not trivially be used), and as part of those changes the appealingly-named `Ref<T>` type in the core module was re-used for this new case. Builtin operations were declared with an explicit `Ref<T>` return type, and parts of the compiler front-end that had previously been blissfully unaware of the AST's `RefType` (and `InOutType`, etc.) had to start accounting for the possibility that an explicit `Ref<T>` would show up. Related changes also introduced a comparable conflation of the (unfortunately-named) `constref` parameter-passing modifier and builtin operations that wanted to return an explicit reference that is read-only. Both use cases were mapped to the core-module `ConstRef<T>` type, which appeared in the AST as an instance of the `ConstRefType` class. The overlapping use of `ConstRef<T>`` is actually significantly more troublesome than the `Ref<T>` case because, despite what its name implies, `constref` was not really supposed to be the read-only analogue of `ref`, but rather it is closer to the "immutable value borrow" analogue to `inout`'s "mutable value borrow." The semantics of a "value borrow" vs. a "memory reference" in Slang have not been very carefully codified, and the conflation around `ConstRef<T>` has contributed to things becoming increasingly muddy in the compiler back-end. Main Changes ============ Core Module ----------- The `Ref<T>` type has been replaced with two distinct types, with one for each use case: * `RefParam<T>` is intended for use when encoding a `ref` parameter in a function type * `ExplicitRef<T>` is intended for use when an operation in a builtin module wants to return a reference The other types used to represent parameter-passing modes (e.g., `InOut<T>`) were renamed to better indicate that their role in defining parameter types (e.g., `InOutParam<T>`). The `ExplicitRef<T>` type was given additional generic parameters for the allowed access and the address space, akin to what `Ptr<T>` now supports. The pointer dereference operator (prefix `*`) in the core module should now properly propagate the access and address space of the pointer over to the reference that gets returned. The two distinct use cases of `ConstRef<T>` were not split in the way as `Ref<T>`, instead the case for the `constref` parameter-passing mode uses `ConstParamRef<T>`, while cases that previously used `ConstRef<T>` to represent a read-only explicit reference instead now use `ExplicitRef<T, Access.Read>`. Prior to this change there were two subscripts declared on pointers: one in the `Ptr` type itself, and another in an `extension` for pointers with `Access.ReadWrite`. The comments on the code seemed to indicate that the catch-all subscript used to only have a `get` accessor, while the `ref` was only available on read-write pointers, but it seems that subsequent changes converted the default subscript to support `ref`. This change eliminates the subscript added via `extension`, since it is redundant. AST and Front-End ================= Similar to the changes in the core module, the AST `RefType` class was split into: * `RefParamType` for the case of encoding `ref` parameters * `ExplicitRefType` for the case where the user meant an explicit reference type All the other classes that represent wrappers for encoding parameter-passing modes (e.g., `OutType`) were similarly renamed (e.g., `OutParamType`). The `ConstRefType` class was simply renamed to `ConstRefParamType`, because any use cases of `ConstRefType` that intended an explicit reference type will now use `ExplicitRefType` with `Acccess.Read`. For convenience, this change includes type aliases to map the old names for these types over to the new ones (e.g., `using OutType = OutParamType`) so that the change doesn't need to affect quite so many lines of code. The `RefType` and `ConstRefType` names are intentionally left undefined, since it woudl be unsafe to assume that existing use sites should default to either of the two possible interpretations. All use cases of `RefType` and `ConstRefType` (and their former shared base class `RefTypeBase`) were audited and updated to refer to either `RefParamType`/`ConstRefParamType` or `ExplicitRefType`, as appropriate (based on whether the context of the code indicated it was working with parameter-passing mode wrapper types, or explicit reference types). In many (many) cases comments were added to the code that was updated (and some unrelated code that needed to be audited along the way) to note cases where there appears to be something fishy going on in the compiler and/or there are obvious opportunities for next-step improvement. The `QualType` constructor used to infer l-value-ness when passed a `RefType` or `ConstRefType`; that code was introduced to support explicit reference types. The code was updated to consult the access argument of an `ExplicitRefType` to try and determine the right l-value-ness to use. There is some ambiguity about what should be done in the case where the value of the generic argument representing the access cannot be statically determined; a better solution may be needed. Many other cases in the front-end that were working with `RefType` and `ConstRefType` for explicit references also need to figure out l-value-ness, and these were changed to rely on the logic already added to `QualType` so that it wouldn't have to be duplicated. It isn't clear if this structure is the best way to tackle the problem, but it seems to at least be an upgrade over the more strictly ad-hoc logic that was in place before. Future Work =========== IR-Level Work ------------- The most obvious next step to take is that the split that was made in the compiler front-end needs to be properly plumbed through all of the back-end. There appears to be a lot of code in the back end of the compiler that has made the same conflation of `ref` parameters and explicit reference types that the front-end did. In practice, any uses of `ExplicitRef<T>` in the front-end should desugar into plain pointer-based code in the IR. Clean Up Parameter-Passing Modes -------------------------------- The code that handles different parameter-passing modes (`ParameterDirection`s) and their wrapper types is somewhat scattered and messy (as found while auditing use cases of `RefType`). A cleanup pass is warranted to ensure that most code only needs to think about `ParameterDirection`s. There should ideally be only a single operation in the front-end that handles determining the `ParameterDirection` of a parameter based on its modifiers. Similarly, there should be one operation to wrap a value type based on a parameter direction, and one operation to derive a `ParameterDirection` from the wrapper type. Ideally, the accessors for `FuncType` should not provide unrestricted access to the potentially-wrapped parameter types, and should instead return some kind of `ParamInfo` struct that encodes both a `ParameterDirection` and the unwrapped `Type` of the parameter. Clean Up `QualType` ------------------- A significant piece of future work that appears required is to drastically clean up and improve the way that `QualType`s are represente and handled in the front-end. There are currently various distinct `bool` flags in `QualType` (some with very unclear meaning) and differnet parts of the codebase consult/modify only subsets of them; a clear enumeration of the "value categories" (to use the C++ terminology) that Slang supports could be quite helpful. Naively, a `QualType` should at least encode the basic information that a `Ptr` type encodes: * A value type * Allowed access (read-only, read-write, etc.) * Address space The main additional thing that a `QualType` needs is a way to distinguish cases where an expression evaluates to: * A reference to a memory location, where all the information from a `Ptr` is relevant * A simple value, such that the access and address space are irrelevant * A reference to an abstract storage location (a `property`, `subscript`, or an implicit conversion that needs to support being an l-value), in which case address space is irrelevant and the "allowed access" basically amounts to a listing of the accessors the storage location supports Eliminate Explicit Reference Types ---------------------------------- Finally, twe should eventually eliminate the `ExplicitRef<T>` type from the core module (and all of the supporting code from the front-end), since the feature is not a good fit for the Slang language. We should find some other way to decorate operations in the builtin module that need to returns a reference rather than a value (note how `ref` accessors already avoided exposing explicit reference types, by design). --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Add RHI Device Caching and Test Prefix Exclusion (#8448)Gangzheng Tong2025-09-22
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | # Add RHI Device Caching and Test Prefix Exclusion ## Summary This PR introduces two key improvements to the Slang test infrastructure: 1. **RHI Device Caching**: Implements device caching to significantly speed up test execution by reusing graphics devices across tests, **RHI Device Caching reduces slang-test execution time from ~15 minutes to ~5 minutes in Windows release builds** 2. **Test Prefix Exclusion**: Adds `-exclude-prefix` option to skip tests matching specified path prefixes ## Changes ### RHI Device Caching - **New `DeviceCache` class** (`slang-test-device-cache.h/cpp`): Thread-safe device cache with LRU eviction (max 10 devices) - **Cache control option**: `-cache-rhi-device` flag in both `slang-test` and `render-test` - Default: **enabled** in slang-test, **disabled** in render-test when run standalone - Automatically skips caching for CUDA devices (due to driver issues) - **Performance benefit**: Eliminates expensive device creation/destruction cycles, especially beneficial for Vulkan on Tegra platforms ### Test Prefix Exclusion - **New `-exclude-prefix <prefix>` option** in slang-test - Allows excluding entire test directories or patterns from execution - Complements existing `-category` and individual test filtering options ### Usage Examples ```bash # Enable device caching (default) slang-test # Disable device caching slang-test -cache-rhi-device false # Exclude tests from specific directories slang-test -exclude-prefix tests/problematic/ slang-test -exclude-prefix tests/slow/ -exclude-prefix tests/experimental/ ``` This change should significantly improve test execution performance, particularly in CI environments with frequent device operations. This is needed for running the GPU test in aarch64, where repeated device creation/destroy is causing driver issues. Needed by: https://github.com/shader-slang/slang/issues/8346 --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Reduce the maxprocesses to 3 for slangpy test (#8502)Gangzheng Tong2025-09-22
|
* Fixed typo in `SemanticsVisitor::_readAggregateValueFromInitializerList` (#8504)Ronan2025-09-22
| | | I think the commit diff speaks for itself.
* Revise SPIRV-Headers update instructions in documentation (#8488)Jay Kwak2025-09-20
| | | | | | It turned out that SPIRV-Headers update has to be done after building SPIRV-Tools Co-authored-by: Gangzheng Tong <tonggangzheng@gmail.com>
* Use LOAD_LIBRARY_SEARCH_DEFAULT_DIRS for LoadLibraryExW (#8491)Gangzheng Tong2025-09-19
| | | | | | | | | | | | | | | | | Before: - Uses `LOAD_LIBRARY_SEARCH_USER_DIRS` in `LoadLibraryExW`, which might cause exception if there is no pathes added by `AddDllDirectory()` After: - Use the composite flag `LOAD_LIBRARY_SEARCH_DEFAULT_DIRS`, which searches for several locations. - Will still search dir added by `AddDllDirectory()`, but avoids empty path seraching if there is no AddDllDirectory() calls. Related to https://github.com/shader-slang/slang/issues/8462 --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Check if any required jobs failed (allow success or skipped) (#8495)Gangzheng Tong2025-09-19
| | | This allows doc only changes to PASS the CI by skipping jobs.
* Fix SPIRV headers precedence (#8481)Jay Kwak2025-09-19
| | | | | | | | | | Remove SYSTEM flag from SPIRV-Headers to fix MacOS header precedence. When the path is registered as SYSTEM, it is used with `-isystem` option not `-I` option and it gets less searching order on MacOS. When spirv.h is installed on the system directory, it will end up using the system installed spirv.h, which is most likely an older version than we should use.
* include limits.h in slang-platform.cpp (#8473)John Zupin2025-09-19
| | | | | | | | fixes https://github.com/shader-slang/slang/issues/8472 Fixes an issue with GCC 9.4.0 on Ubuntu 20.04, it will throw an error about PATH_MAX not being declared. Co-authored-by: Mukund Keshava <mkeshava@nvidia.com>
* Fix DebugCompilationUnit to reference main shader file instead of header ↵Lujin Wang2025-09-18
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | files (#7957) This PR implements the requested fix for issue #7923 where DebugCompilationUnit incorrectly referenced header files instead of the main shader file. ## Summary - Modified IRDebugSource to include isIncludedFile flag as third operand - Updated emitDebugSource function to accept and pass the included file flag - Updated call sites to use source->isIncludedFile() from SourceFile class - Modified SPIR-V emission to only create DebugCompilationUnit for non-included files ## Test Results The fix has been verified with the provided reproducer code. The SPIR-V output now correctly shows DebugCompilationUnit referencing the main shader file instead of header files. Generated with [Claude Code](https://claude.ai/code) --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: Lujin Wang <lujinwangnv@users.noreply.github.com> Co-authored-by: Claude Code <claude@anthropic.com> Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Fix CUDA global variable initialization with constructor calls (#8340)Harsh Aggarwal (NVIDIA)2025-09-18
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Fix CUDA global variable initialization with constructor calls Resolves CUDA compilation failure where global variables with struct constructor initialization generated illegal `__device__` variable runtime initialization. **Problem:** ```cuda // Generated invalid CUDA code: __device__ static const Stuff_0 gStuff_0 = Stuff_x24init_0(args...); // Error: "dynamic initialization is not supported for a __device__ variable" Root Cause Discovered: Through extensive debugging, found that moveGlobalVarInitializationToEntryPoints pass only handled kIROp_GlobalVar instructions, but global constants with constructor calls appeared as kIROp_Call instructions at module scope. Solution: 1. IR Pipeline Fix: Extended moveGlobalVarInitializationToEntryPoints to detect and transform module-level constructor calls into proper global variables with entry-point initialization 2. Field Access Fix: Enhanced kIROp_FieldExtract logic to emit correct -> syntax for pointer types and address-of operations 3. Constructor Emission: Added CUDA-specific handling for constructor calls Architecture: - Transforms let %gStuff = call %Constructor(...) into kernel context initialization - Moves runtime initialization from global scope to entry-point execution - Follows CUDA best practices for global state management Files: - source/slang/slang-ir-explicit-global-init.cpp: Extended IR transformation pass - source/slang/slang-emit-c-like.cpp: Enhanced field access and foldable value logic - source/slang/slang-emit-cuda.cpp: Added CUDA-specific field extraction handling Result: // Now generates proper CUDA code: struct KernelContext_0 { Stuff_0 gStuff_1; }; // Runtime initialization in entry point: kernelContext_1.gStuff_1 = constructor_call(); Fixes: tests/compute/type-legalize-global-with-init.slang
* Add WASM FS module support for slang-playground (#8459)Gangzheng Tong2025-09-17
| | | | | | | | | | | | | Add WASM FS module support for slang-playground This change adds the necessary Emscripten build flags to export the FileSystem (FS) module interface in the slang-wasm build: - Adds -sMODULARIZE=1 to enable modular builds - Adds -sEXPORTED_RUNTIME_METHODS=['FS'] to export the FS interface These changes are required to support the slang-playground. The existing flags are also reformatted for better readability. Related to https://github.com/shader-slang/slang-playground/issues/170
* Set default mimalloc for spirv-tools but allow user to override (#8460)Gangzheng Tong2025-09-17
| | | | | | | | | | | | | | | | | | | | | | | | This PR resolves a configuration conflict where the SLANG_ENABLE_SPIRV_TOOLS_MIMALLOC option was defined twice with different defaults, causing user settings to be overridden. Changes: - Removed duplicate definition in external/CMakeLists.txt that only applied to WIN32 - Define default in main CMakeLists.txt - ON by default for Windows - OFF by default for other platforms (still need to resolve crashes so set it to `OFF` for now) Preserved user configuration - explicit -DSLANG_ENABLE_SPIRV_TOOLS_MIMALLOC=value settings are now properly respected Related to https://github.com/shader-slang/slang/issues/8158 --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Diagnostic for metal ref mesh output assignment (#8365)James Helferty (NVIDIA)2025-09-17
| | | | | | | When slang detects assignment to a mesh output reference on metal, generate a diagnostic message. (Metal mesh shader outputs must be assigned via 'set' instead of 'ref'.) Fixes #7498
* Use self-hosted runner for Windows release build (#8470)Gangzheng Tong2025-09-17
| | | | | | | | | | - Remove hard-coded Win SDK version. - Using self-hosted machine for building release package in Windows. - Remove the steps from release.yml that have been done in common steup. --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Fix LSS intrinsics for hit objects in ray tracing tests (#8469)Harsh Aggarwal (NVIDIA)2025-09-17
| | | | | | | | | | | | Enable GetLssPositionsAndRadii() call in rayGenLssIntrinsicsHitObject shader that was previously commented out. This fixes the failing ray-tracing-lss-intrinsics-hit-object test which was returning all zero values for LSS position and radius data. The hit object LSS intrinsics are now working correctly in D3D12 backend, returning proper endcap positions and radii values as expected by the test. All 27 test assertions now pass successfully. Fixes #8128
* Added __magic_enum (#8436)Ronan2025-09-17
| | | | | | | | | | | | | | | | | | | | | Fixes #8406 (and #8410). `AddressSpace`, `MemoryScope` and `AccessQualifier` are no longer `BaseType`. I added a new `__magic_enum` (very similar to `__magic_type`) syntax to be able to easily create values or these enums from the compiler. (I don't know if it was the right way to do it, but it works and the changes are small enough?). I had a weird bug: `tests/language-feature/capability/address-of.slang` was failing in `IRBuilder::_findOrEmitConstant(IRConstant& keyInst)`. When needing a new `u64(0)`, it did not find it in the `ConstantMap` first, but then failed to add it right after because it already existed in the map! But this was triggered by `IRPtrType* IRBuilder::getPtrType(IROp op, IRType* valueType, AccessQualifier accessQualifier, AddressSpace addressSpace)`, which is a strange coincidence... but I could not find the issue in what I did. I ended up bumping unordered_dense, and it solved the issue (so there was a bug in there).
* Change CI to use locally installed sdk (#8450)Mukund Keshava2025-09-17
|
* Diagnose error when the function args can't satisfy constexpr parameter ↵Gangzheng Tong2025-09-16
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | requirements (#7269) ## Summary This PR enhances constexpr validation by adding proper error checking when function arguments cannot satisfy constexpr parameter requirements, addressing issue #6370. ## Problem Previously, when a function declared constexpr parameters, the compiler would attempt to propagate constexpr-ness to the call site arguments, but there was insufficient validation and error reporting when this propagation failed. This could lead silent failures where constexpr requirements weren't properly enforced ## Solution This PR adds checks that: 1. **Validates constexpr arguments**: When a function parameter is marked as `constexpr`, the compiler now explicitly checks that the corresponding argument can be marked as `constexpr` 2. **Issues clear compilation errors**: added `Diagnostics::argIsNotConstexpr`) 3. **Handles both call scenarios**: The validation works for both: - Direct function calls with IR-level function definitions - Calls to function from external modules Fixes #6370 --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Enable static mimalloc in Spriv-Tools (#8419)Gangzheng Tong2025-09-16
| | | | | | | | | | | | | | | | | | | | | | * Added optional mimalloc integration for SPIRV-Tools to improve compilation performance * Cloning `mimalloc` repo to external/ on demand during cmake config * Enabled by default on Windows, configurable via `SLANG_ENABLE_SPIRV_TOOLS_MIMALLOC` CMake option, to enable static mimalloc build in spirv-tools. * There are some crashes in Linux and Mac with static mimalloc enabled, likely due to the system malloc and mimalloc mixed usage. This might be expected as mimalloc in spirv-tools is not extensively tested according to https://github.com/KhronosGroup/SPIRV-Tools?tab=readme-ov-file#dependency-on-mimalloc. So by default only Windows has this enabled. Close: https://github.com/shader-slang/slang/issues/8158 --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Fix segfault in SPIR-V header processing in SpirvInstructionHelper (#8428)Gangzheng Tong2025-09-10
| | | | | | | | | | | | | | | | | | | | The `SpirvInstructionHelper::loadBlob()` method could segfault when calling `m_headerWords.addRange()` if the SPIR-V blob contained insufficient data for the required 5-word header. To reproduce, run ``` ./build/Debug/bin/slangc.exe tests/modules/environment.slang -o tests/modules/environment.slang-module -target spirv -separate-debug-info (0): error 57004: output SPIR-V contains no exported symbols. Please make sure to specify at least one entrypoint. Segmentation fault ``` The error is expected, but the `Segmentation fault` is not. This PR adds the check to ensure the SPIR-V blob has at least `SPV_INDEX_INSTRUCTION_START * sizeof(SpvWord)` bytes (20 bytes minimum) before attempting to process the header words. Related to: https://github.com/shader-slang/slang/issues/7547
* Add FindModifier for Declarations (#8308)Xuanda Yang2025-09-10
| | | | | | Add `findModifier` for `DeclReflection` so pattern like `extern struct foo;` can be properly reflected. Closes #8009
* Fix crash when compiling specialized generic entrypoint containing a static ↵Yong He2025-09-10
| | | | | | | | | | | | | | | | | | | | | const decl. (#8392) Closes #8184. We fixed three issues with this regression test: 1. After generating IR for a `SpecializeComponentType`, we should also strip the frontend decorations from the IR so there is no HighLevelDeclDecoration that will go into the backend. 2. When lowering a static const inside a generic function, we should not give the static const a linkage, because it won't such constant will not appear in global scope. Trying to give it a linkage decoration will lead to the parent generic (for the function) to have two duplicate Export/Import decorations with different mangle names, and confuses the linker. 3. Make sure internal exceptions does not leak through `IComponentType::getEntryPointCode`/`getTargetCode`.
* Fix pointers and C-like layout in varying parameters (#8425)Julius Ikkala2025-09-10
| | | | | | | | | | | | | Closes #8409, but ended up being more about fixing another bug. While the issue itself seems to only be a simple typo fix (see second commit in this PR), I found out during writing a test that pointers never got correct locations regardless of layout. Their locations were always assigned to zero due to lacking a resource usage entry in `TypeLayout`. They were also missing the `Flat` decoration, so I went ahead and added that too. I can split this up into two separate PRs if that's preferred; both aspects just share a test right now and fix a similar-looking issue in the resulting SPIR-V.
* Squash warnings on gcc 14 (#8377)Ellie Hermaszewska2025-09-10
|
* CUDA: Fix compiler crash with unsized array field - nonuniformres-as-… (#8380)Harsh Aggarwal (NVIDIA)2025-09-10
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | …function-parameter.slang #8315 Root Cause: CUDA compilation crashed with `assert failure: !seenFinalUnsizedArrayField` because unsized arrays like `RWStructuredBuffer<uint> globalBuffer[]` were not the final field in generated parameter structs, violating the layout constraint in slang-ir-layout.cpp. Fix: Extended `collectGlobalUniformParameters` to automatically reorder struct fields for CUDA targets - regular fields first, unsized arrays last. Other targets preserve original order. Impact: - Enables CUDA support for nonuniform resource indexing as function parameters - Zero impact on existing GLSL/HLSL/SPIRV targets - Automatic handling - no manual parameter reordering required Files: slang-emit.cpp, slang-ir-collect-global-uniforms.cpp/.h, test file --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com> Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
* Check if debugVar for is debuggable types in the legalization pass (#8326)Gangzheng Tong2025-09-10
| | | | | | | | | | | | | | | | | ## Problem When generic functions with debug variables were specialized with concrete types containing non-debuggable fields (e.g., `StructuredBuffer`), the IR cloning process would create invalid `DebugVar` instructions without checking if the substituted types remained debuggable. ## Solution This fix adds a defensive check in the legalization pass that removes the debugVar created for the non-debuggable types. --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Fix language server auto-complete regression in debug build. (#8416)Yong He2025-09-10
| | | | | | | | | | | | | | | | | | | | | | | | | | | Fixes this regression: ```slang struct MyType { // Regression Condition 1: there must be more than one member in the lookup scope. float v; int getSum() { return 0; } } void m(MyType t) { // Regression condition 2: the completion must be in an init expression. // Regression condition 3: none of the candidate members can coerce to the expected type. // Regression behavior: no completion candidates are shown, because // SemanticsVisitor::resolveOverloadedLookup throws an error when there are 0 applicable candidates // after type coercion filtering. Texture2D x = t.; // completion request after . here } ``` The root cause is that we shouldn't be applying candidate filtering on the candidate list when in completion checking mode. Closes #8417.
* Update SPIRV-{Tools,Headers} (#8415)Jay Kwak2025-09-10
| | | | | | | | | This updates SPIRV submodules. Unfortunately we cannot use the latest because one of PRs in SPIRV-Tools causes VVL errors on our slang-tests. The issue is tracked in - https://github.com/KhronosGroup/SPIRV-Tools/issues/6316 - https://github.com/shader-slang/slang/issues/8414 This PR also updates/simplifies the document of updating SPIRV process.
* Send notification only on scheduled CTS runs (#8418)Gangzheng Tong2025-09-09
|
* Enable ccache in slang bulid for CTS test (#8413)Gangzheng Tong2025-09-10
| | | | - Enable ccache in slang bulid for CTS test - Allow manual CI dispatch
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-03 00:14:06 +0000