| Commit message (Collapse) | Author | Age |
|---|
| ... | |
| |
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
| |
Related to
- https://github.com/shader-slang/slang/issues/8519
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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
|
| |
|
|
|
|
|
|
|
|
|
| |
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.
|
| |
|
|
|
|
|
|
|
|
|
|
| |
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.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
| |
I think the commit diff speaks for itself.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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
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
|
| |
|
|
|
|
|
| |
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
|
| |
|
|
|
|
|
|
|
|
|
|
| |
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
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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 `DeclReflection` so pattern like `extern struct
foo;` can be properly reflected.
Closes #8009
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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`.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
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.
|
| | |
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
…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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
## 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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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.
|
| |
|
|
|
|
|
| |
macos 15.6 includes python 3.9.6 with Xcode, which doesn't understand
match/case. Changing it to to the less spiffy if/elif.
Co-authored-by: Yong He <yonghe@outlook.com>
Co-authored-by: Sam Estep <sam@samestep.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Fixes #8396 by not emitting the `ArrayStride` when it would've been
zero. The problem is caused by #7848, more details in the issue
description.
I checked that with equivalent GLSL code, glslangValidator does not emit
any `ArrayStride`. I assume that the addition of `storageClass ==
SpvStorageClassStorageBuffer` to line 1848 is not a mistake. If it is,
that addition could also be simply reverted to fix this issue, I tested
that option as well. With these changes, Slang tests work locally on my
PC again.
Related to this; it'd be nice to have GPUs from multiple vendors in the
CI to avoid this kind of thing happening again. Or even just llvmpipe;
that doesn't even require a GPU and would add at least one more driver
to test with.
|
| |
|
|
|
|
|
|
|
| |
This change relaxes a previous restriction on link-time types and
constants, so that we now allow them to be used to define shader
parameters.
Doing so will result in a parameter layout that is incomplete prior to
linking. The PR added a test to call the reflection API on a fully
linked program and ensure that we can report correct binding info.
|
| |
|
|
|
| |
The code int x4 = 0xFFFFFFFFFFFFFFFF previously did not produce a
warning due to the value being too large for the type. This patch now
checks for this and similar issues during parsing.
|
| |
|
| |
Update intrinsics signature as per the nvapi header
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Enable CUDA support for additional HLSL intrinsic tests by implementing
missing functionality and fixing compiler bugs affecting CUDA targets.
- Fix critical bug in InterlockedCompareStore64 where division used /4
instead of /8 for 64-bit types, causing incorrect memory addressing for
all signed int 64_t atomics
- Add signed int64_t atomic wrappers (atomicExch, atomicCAS) to CUDA
prelu de that properly cast to/from unsigned types as required by CUDA's
atomic API
- Enable tests: atomic-intrinsics-64bit.slang
- Implement CUDA support for QuadAny and QuadAll operations using warp
shu ffle primitives (__shfl_sync with quad-level lane masking)
- Add CUDA to quad_control capability definition in
slang-capabilities.capdef
- Add _slang_quadAny/_slang_quadAll helper functions to CUDA prelude
- Enable tests: quad-control-comp-functionality.slang,
subgroup-quad.slang
---------
Co-authored-by: szihs <675653+szihs@users.noreply.github.com>
|
| |
|
| |
closes https://github.com/shader-slang/slang/issues/3313
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
This commit is to emit the debug-info for the entry point parameters.
Two things are implemented/fixed in this PR:
- We were not emitting the `DebugVar` and `DebugValue` at the IR
lowering level when the type of the entry point parameter is `ConstRef`.
This commit handles the `ConstRef` case in a same way that the other
types are handled so that `DebugVar` and `DebugValues` are properly
emitted at the IR lowering level.
- Two types for Geometry shaders were incorrectly treated as not valid
types for the DebugInfo. They are `InputPatch` and `OutputPatch`. This
commit handles them as valid types for DebugInfo.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
| |
Update docs/shader-execution-reordering.md with additional intrinsics
Add correct capability `LoadLocalRootTableConstant`
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Emits the appropriate OpCapability for 8- and 16-bit type usage:
- UniformAndStorageBuffer8BitAccess: for 16-bit types in
SpvStorageClassUniform and SpvStorageClassStorageBuffer
- UniformAndStorageBuffer16BitAccess: for 16-bit types in
SpvStorageClassUniform and SpvStorageClassStorageBuffer
- StoragePushConstant8: for 8-bit types in SpvStorageClassPushConstant
- StoragePushConstant16: for 16-bit types in SpvStorageClassPushConstant
- StorageInputOutput16: for 16-bit types in SpvStorageClassInput and
SpvStorageClassOutput
Generated with Claude Code, with revisions.
Fixes #7879.
---------
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: James Helferty (NVIDIA) <jhelferty-nv@users.noreply.github.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#8223)
The Metal backend was generating incorrect type names for 8-bit vector
types, causing compilation failures when targeting Metal. According to
the Metal specification, 8-bit vector types should be named `charN` and
`ucharN` (e.g., `char2`, `uchar3`) rather than `int8_tN` and `uint8_tN`.
## Problem
When compiling Slang code with 8-bit vector types for Metal, the
compiler would emit:
```metal
uint8_t2 _S8 = uint8_t2(uint8_t(0U), uint8_t(16U));
int8_t3 _S9 = int8_t3(int8_t(0), int8_t(16), int8_t(48));
```
But the Metal compiler expects:
```metal
uchar2 _S8 = uchar2(uint8_t(0U), uint8_t(16U));
char3 _S9 = char3(int8_t(0), int8_t(16), int8_t(48));
```
This caused errors like:
```
error: unknown type name 'uint8_t2'; did you mean 'uint8_t'?
```
## Solution
Modified `MetalSourceEmitter::emitSimpleTypeImpl()` to emit the correct
Metal-specific type names for 8-bit types:
- `kIROp_Int8Type` now emits `char` instead of `int8_t`
- `kIROp_UInt8Type` now emits `uchar` instead of `uint8_t`
This change only affects the Metal backend and ensures that vector types
like `int8_t2`, `uint8_t3`, etc. are correctly emitted as `char2`,
`uchar3`, etc.
## Testing
- Added a new test case `tests/metal/8bit-vector-types.slang` to verify
the fix
- Re-enabled the previously disabled Metal test in
`tests/hlsl-intrinsic/countbits8.slang`
- Updated `tests/metal/byte-address-buffer.slang` to expect the correct
type names
- Verified that existing Metal tests continue to pass
Fixes #8211.
<!-- 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: bmillsNV <163073245+bmillsNV@users.noreply.github.com>
|
| |
|
|
|
|
| |
Fixes: #8018
Changes:
* Do not emit true for `shouldEmitSPIRVDirectly` with a GLSL target
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
close #8068.
Currently the AutoDiff aggressively scan every IR inst in searching the
differentiable IR. This is not efficient and could have bug, details in
https://github.com/shader-slang/slang/issues/8068#issuecomment-3214856668.
This PR change the behavior. It will do a initial filter to only gather
the global differentiable IRs and IRFunc and IRGeneric as well. For
IRGeneric, we will pick it only when it's used in other generic function
(it's only useful when dealing with dynamic dispatch).
Then we will start searching reachable insts from this IR list by using
the same method as before.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#7929)
Fixes the Slang compiler internal error "subscript had no getter" when
reading from mesh shader output index arrays (e.g., `triangles[0].x`).
## Problem
The `OutputIndices` struct was missing a `ref` accessor in its
`__subscript` implementation, causing the compiler to fail when trying
to materialize subscript expressions as r-values.
## Solution
Added the missing `ref` accessor to `OutputIndices.__subscript` using
the `kIROp_MeshOutputRef` intrinsic operation, matching the pattern used
in `OutputVertices` and `OutputPrimitives`.
## Files Changed
- `source/slang/core.meta.slang` - Added missing `ref` accessor
- `tests/bugs/gh-7925.slang` - Test case to reproduce and verify the fix
Fixes #7925
Generated with [Claude Code](https://claude.ai/code)
---------
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Lujin Wang <lujinwangnv@users.noreply.github.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
| |
`emitReflectionVarLayoutJSON` will output the `userAttribs` section
twice as it gets output by `emitReflectionModifierInfoJSON` first before
being output again by a direct call to `emitUserAttributes`.
It seems the answer here is to just remove the extra explicit call to
`emitUserAttributes` and rely on the call in
`emitReflectionModifierInfoJSON`?
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Closes #8112. ~~The issue asks for a "C layout", but in this PR I use
the term "CPU layout" because this naming was pre-existing in the
codebase as `kCPULayoutRulesImpl_`. The primary purpose of this layout
is to match CPU-side struct definitions with the shader side. I'm open
to better naming suggestions, though.~~
Edit: switched back to using `CDataLayout` & `-fvk-use-c-layout`, as the
CPU target depends on the object layout rules of existing CPU layout
rules, but they're incompatible with actual shaders. So a new
`kCLayoutRulesImpl_` was needed anyway.
---------
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
-Adds semantic SV_VulkanSamplePosition that emits corresponding
gl_SamplePosition and SpvBuiltinSamplePosition
-Adds gl_SamplePosition property to glsl.meta.slang
-Adds SPIRV and GLSL tests for the semantic and property
-Plan is to later implement SV_SamplePosition that follows HLSL range of
-0.5 to +0.5,
and emits GetRenderTargetSamplePosition(SV_SampleIndex) which needs more
complicated IR manipulation for HLSL and Metal
Fixes #7906
---------
Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
| |
Fixes #8185. The previous implementation is incorrect and basically only
works in the `x = 0` case. `delta` was the smallest possible positive
value representable as a float, but that's below the rounding error of
addition with almost all reasonably sized floats.
This fixed implementation is based on bit twiddling instead. I've
checked the float case against the C++ `nextafterf` with both a -inf ->
inf and inf -> -inf sweep, in addition to the test included in this PR.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
## Summary
- Add Metal platform support for `WaveGetActiveMask()` and
`WaveActiveCountBits()` wave intrinsics
- Update capability requirements to include Metal platform for subgroup
ballot operations
- Implement Metal-specific intrinsic assembly using `simd_ballot()` and
`simd_vote` APIs
## Changes
- **source/slang/hlsl.meta.slang**:
- Add Metal target case for `WaveGetActiveMask()` using
`simd_ballot(true)`
- Update capability requirements from `cuda_glsl_hlsl_spirv` to
`cuda_glsl_hlsl_metal_spirv` for wave ballot functions
- **source/slang/slang-capabilities.capdef**:
- Add `metal` to `subgroup_ballot_activemask` capability alias
|
| |
|
|
|
|
|
|
|
| |
Enable CUDA support for batch 3 tests
- Enhanced wave operations with exclusive support
- Added proper identity values for min/max operations
- Fixed intrinsic name mapping issues
- Updated test configurations
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
|
|
|
|
|
| |
In Metal, if `ParameterBlock` contains `DescriptorHandle` directly, it
would be emitted as DescriptorHandle literal, which is not valid Metal
code,
This fix adds a case for `kIROp_DescriptorHandleType` and directs it to
the Parent's `emitType` function to handle it.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
This is a followup on #7828 to fix bugs that were causing CodeLLDB to
give wrong values and hang (see vadimcn/codelldb#1302) because I didn't
realize that these data formatters can be passed _either_ a value of a
given type _or_ a pointer to a value of that type, and need to handle
both cases. I also introduced loop bounds to prevent hangs in the case
where these synthetic values are constructed for things like
uninitialized variables.
From looking at the preexisting data formatters from #4272 in
`source/core/core_lldb.py`, it seems like they _technically_ have
similar bugs to this, but since those types are simpler, it's unclear to
me whether that can actually manifest in meaningful ways like these bugs
in `source/slang/slang_lldb.py` were doing.
Anyways, to test this, put a breakpoint here:
https://github.com/shader-slang/slang/blob/6d399804a353154259cf4410940f144db8f9b5cf/source/slang/slang-emit-cpp.cpp#L1733
And use this `.vscode/launch.json` for CodeLLDB:
```json
{
"version": "0.2.0",
"configurations": [
{
"name": "LLDB",
"preLaunchTask": "Debug build",
"type": "lldb",
"request": "launch",
"initCommands": ["command source .lldbinit"],
"program": "build/Debug/bin/slangc",
"args": [
"tests/cpu-program/cpu-hello-world-test.slang",
"-target",
"executable",
"-o",
"hello"
]
}
]
}
```
Before this PR, the `inst` variable will display in the debug pane as
`{kIROp_StringLit 0x00007fffffff5f68}`, which is the wrong pointer
value. You can also check this by running `p inst` in the Debug Console,
which will print this:
```
(Slang::IRInst *) 0x000055555fdac3b8 {kIROp_StringLit 0x00007fffffff5f68}
```
In contrast, running `p *inst` prints the correct pointer value:
```
(Slang::IRInst) {kIROp_StringLit 0x000055555fdac3b8} {
[op] = kIROp_StringLit
[UID] = 76
[type] = 0x000055555fdac348 {kIROp_StringType None}
[decorations/children] = {}
[parent] = 0x000055555fdac2d0 {kIROp_ModuleInst None}
[uses] = 0x000055555fdadf18 {kIROp_StringLit 0x000055555fdac3b8}
}
```
But as you can see, in that case the synthetic `[value]` child is
completely missing.
Then if you try to expand `inst` in the debug pane, CodeLLDB will hang
(or at least it does when I try this).
After this PR, the hex integer for the pointer is always consistent, and
CodeLLDB does not hang in the debug pane when you expand `inst`, and
shows the correct `[value]` child just like when running `v *inst`.
As an aside: after this PR, the `[value]` child is still missing when
specifically running `p *inst` in the Debug Console. It _is_ possible to
fix this:
```diff
diff --git a/source/slang/slang_lldb.py b/source/slang/slang_lldb.py
index 23905d8c5..d2b3a4da9 100644
--- a/source/slang/slang_lldb.py
+++ b/source/slang/slang_lldb.py
@@ -93,13 +93,11 @@ class IRInst_synthetic(lldb.SBSyntheticValueProvider):
value: list[tuple[str, lldb.SBValue]] = []
match op.value:
case "kIROp_StringLit":
- string_lit_t = target.FindFirstType("Slang::IRStringLit")
- string_lit = self.valobj.Cast(string_lit_t)
+ string_lit = self.valobj.EvaluateExpression("(Slang::IRStringLit*)this")
val = string_lit.GetChildMemberWithName("value")
value = [("[value]", val.GetChildMemberWithName("stringVal"))]
case "kIROp_IntLit":
- int_lit_t = target.FindFirstType("Slang::IRIntLit")
- int_lit = self.valobj.Cast(int_lit_t)
+ int_lit = self.valobj.EvaluateExpression("(Slang::IRIntLit*)this")
val = int_lit.GetChildMemberWithName("value")
value = [("[value]", val.GetChildMemberWithName("intVal"))]
diff --git a/typings/lldb.pyi b/typings/lldb.pyi
index 2672ba244..3a08e9141 100644
--- a/typings/lldb.pyi
+++ b/typings/lldb.pyi
@@ -496,7 +496,7 @@ class SBValue:
def Persist(self): ...
def GetDescription(self, description): ...
def GetExpressionPath(self, *args): ...
- def EvaluateExpression(self, *args): ...
+ def EvaluateExpression(self, expr: str) -> SBValue: ...
def Watch(self, *args): ...
def WatchPointee(self, resolve_location, read, write, error): ...
def GetVTable(self): ...
```
However, that makes the debugger run _significantly_ slower, so I'm
choosing not do do it here.
---------
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
The documentation added by #6844 included instructions to make sure that
the Fiddle `#include` in a file comes after all the other `#include`s,
but it's easy to accidentally violate this via `clang-format`, as
happened for `source/slang/slang-ast-modifier.h` in #7559. This PR
guards against this sort of violation by separating all Fiddle
`#include`s from other `#include`s via a blank line followed by a `//`
line (as we already do in most cases), and also adds a sentence about
this in `tools/slang-fiddle/README.md`.
As a bonus, I also enabled Markdown syntax highlighting for all the code
blocks in that doc file.
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
|
|
| |
Close #8054.
For detailed root cause is at:
https://github.com/shader-slang/slang/issues/8054#issuecomment-3189579508
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Close #8090.
When we do type coerce, we use a cache to store the conversion cost
of different type. The key of the cache is defined by
struct BasicTypeKey
{
uint32_t baseType : 8;
uint32_t dim1 : 4;
uint32_t dim2 : 4;
...
}
where dim1 and dim2 is used for dimension of vector and matrix.
However the dim is only 4 bits, so `vector<int, 16>` will have the same
key as `int`, which is wrong.
Fix the issue by extending it to 8 bit.
Also to make the hash key still within 32 bits, we adjust baseType to 5 bits,
and knownConstantBitCount to 6 bits.
---------
Co-authored-by: kaizhangNV <kazhang@nvidia.com>
|
| |
|
|
|
|
|
|
| |
Add helper functions to create ISlangBlob and load module data from
source.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
