| Commit message (Collapse) | Author | Age |
|---|
| |
|
|
| |
Random drive-by test fix, this was reading past the end of the buffer
but usually succeeded because the expected result is 0.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
This PR implements `Access.Immutable` to allow pointers to immutable
data.
The new type `ImmutablePtr<T>` is defined as an alias of `Ptr<T,
Address.Immutable>`.
By forming a immutable pointer, the programmer is conveying to the
compiler that the data at the pointer address will never change during
the execution of the current program. Therefore loads from immutable
pointers can be deduplicated by the compiler, and will translate to
`__ldg` when generating code for CUDA.
The SPIRV backend is not changed in this PR, since the current SPIRV
spec makes it very difficult to specify loads from immutable address
without generating tons of wrappers and boilerplate type declarations.
We would like to see the spec evolved a bit to around its support of
`NonWritable` physical storage pointers or immutable loads before we
attempt to express such immutability in SPIRV. For now we simply emit
ordinary pointers and loads when generating spirv.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
| |
Closes https://github.com/shader-slang/slang/issues/8154 However there
is further design work to do on implementing the "NonAddressableType"
suggestion
|
| | |
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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.)
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
`Optional<T>`. (#8677)
Closes #8673.
The issue is that we use the RTTI field of an existential to check if it
is null. We have the logic to help the user to fill in a non-zero value
for the RTTI field when such an object is filled from the host. However,
when there is slang code creating an existential value, we still have
old logic in the compiler that just fills in 0 for the RTTI field,
causing an `Optional<IFoo>.hasValue` to always return false in such
cases.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Fixes: https://github.com/shader-slang/slang/issues/7634
Duplicate of PR https://github.com/shader-slang/slang/pull/8052
Primary Changes:
* Added `storeCoherent` and `loadCoherent` for coherent load/store via
pointers. This is backed by `IRMemoryScopeAttr` which is an `IRAttr`
attached to `IRLoad` and `IRStore`
* Logic in `source\slang\slang-emit-spirv.cpp` for load/store emitting
has been reworked to be less messy and more maintainable
* Add to `hlsl.meta.slang` coop vector and coop matrix coherent
load/store operations
Secondary Changes:
* Added a missing load/store test for coop matrix:
`tests\cooperative-matrix\load-store-pointer.slang`
---------
Co-authored-by: ArielG-NV <aglasroth@nvidia.com>
Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Nathan V. Morrical <natemorrical@gmail.com>
|
| |
|
| |
Closes https://github.com/shader-slang/slang/issues/8378
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Add built-in type aliases for DepthTexture* and unify Sampler*Shadow
Add the following type aliases:
- DepthTexture1D, DepthTexture1DArray
- DepthTexture2D, DepthTexture2DArray
- DepthTexture2DMS, DepthTexture2DMSArray
- DepthTexture3D
- DepthTextureCube, DepthTextureCubeArray
These match with the type aliases for non-depth textures.
Also, unify the Sampler*Shadow type aliases with DepthTexture*
ones. This adds the following:
- Sampler2DMSShadow
- Sampler2DMSArrayShadow
and removes the Sampler3DArrayShadow type alias. As a side-effect, the
descriptions of Sampler*ArrayShadow type aliases are fixed
("texture-sampler for shadow" ==> "texture-sampler array for shadow").
Update the slang tests to use the newly introduced type aliases instead
of
the custom type aliases that use _Texture<> directly.
Add DepthTexture testing in
hlsl-intrinsic/texture/texture-intrinsics. Do this by extracting the
test logic of computeMain() in a separate function and parametrize it
for non-depth/depth texture types. This adds basic coverage for the
following types:
- DepthTexture1D
- DepthTexture2D
- DepthTexture3D
- DepthTextureCube
- DepthTexture1DArray
- DepthTexture2DArray
- DepthTextureCubeArray
Issue #6166
Issue #8503
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Closes #8664.
The problem is that when there is an `in` parameter, Slang will create a
local variable to proxy the parameter, copy the value of the parameter
into the proxy variable, and replace all uses of the parameter in the
function body to use the proxy variable instead. This way all writes to
the parameter become writes to the proxy variable.
However, when there is debug info enabled, we are also going to create a
"debugVariable" corresponding to the parameter, but this debugVariable
isn't updated when the proxy variable is updated. The fix is to map the
proxy var instead of the original param to the debug var during the
`insertDebugValueStore` pass, so that any changes to the proxy var will
result in additional stores being inserted to the debug var.
Allowing function body to modify an `in` parameter is a bad legacy
behavior we inherited from HLSL that we should really be moving away
from. I would like us to completely treat an `in` parameter as immutable
by default in the next language version (Slang 2026), and make it an
error if the user tries to do so. This will allow us to generate much
cleaner code and in many cases would help with performance.
|
| |
|
|
|
|
|
|
|
|
|
|
| |
Fix a regression on metal test.
In `lowerBufferElementTypeToStorageType` pass, not only we want to defer
an argument that is `CastStorageToLogical` to the callee, but also apply
the same defer logic to `CastStorageToLogicalDeref` as well.
Because `CastStorageToLogicalDeref` will appear as argumnet if
`lowerBufferElementTypeToStorageType` is run before we apply the
`in->borrow` transformation pass, which is the case for metal parameter
block legalization.
|
| |
|
|
|
|
|
|
| |
happen. (#8653)
This allows us to further cleanup unnecessary copies in the target code
we generate.
Part of effort of #8652.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#8647)
Fixes #8545 where Slang generates SPIR-V with DerivativeGroupQuadsKHR
execution mode but doesn't validate workgroup sizes when texture
sampling triggers automatic derivative computation.
**Root Cause**: Validation code was looking for IRNumThreadsDecoration
on the wrong IR node
**Fix**: One-line change in slang-emit-spirv.cpp to search decoration on
entryPoint instead of entryPointDecor
**Tests**: Added regression tests for both quad and linear derivative
group validation
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: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Fix for a linked list usage bug; avoids dropping any modifiers when
moving type modifiers from a linked list of modifiers into their own
linked list. Since this change results in no_diff modifiers to
traditional functions ending up on the return type instead of the
function (due to the order they're parsed in), we duplicate the no_diff
modifier onto the function declaration after the fact.
Includes a test for the original issue. The no_diff redistribution case
is covered by a slangpy device test case.
Fixes #8332
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
| |
`MakeVector/MatrixFromScalar` (#8512)
- Allows using `Vector/Matrix` type with yet unresolved dimensions
- Simpler implementation and in-line with default `Array`
- Added `test/bugs/gh-8512.slang`
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
We insert field initialization logic at the beginning of every ctor in
`synthesizeCtorBody`, but then immediately inserts another round of
initialization again for explicit ctors in `maybeInsertDefaultInitExpr`,
both called from `SemanticsDeclBodyVisitor::visitAggTypeDecl` right next
to each other.
The fix is to remove `maybeInsertDefaultInitExpr`.
This change also enhances the address aliasing analysis, so that for the
following case:
```
this->member1 = 0;
this->member2 = 0;
this->member1 = param;
```
We can still remove the first assignment to `this->member1` despite
seeing `this->member2=0`, since it is easy to know that `this->member2`
cannot alias with `this->member1`.
Closes #8600.
|
| |
|
|
|
|
|
|
| |
WGSL doesn't allow subgroup related functions in a branching. It must be
used in a uniform flow. This commit disables a test for such case.
Note that the test was supposed to be disabled on the previous PR, but
it was mistakenly not disabled.
- #8386
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#8603)
This change achieves link-time type resolution with a different
mechanism.
For `extern struct Foo : IFoo = FooImpl;`,
instead of synthesizing a wrapper type `Foo` that has a `FooImpl inner`
field and dispatches all interface method calls to `inner.method()`,
this PR completely removes this synthesis step, and instead just lower
such `extern`/`export` types as `IRSymbolAlias` instructions that is
just a reference to the type being wrapped.
Then we extend the linker logic to clone the referenced symbol instead
of the SymbolAlias insts itself during linking.
By doing so, we greatly simply the logic need to support link-time
types, and achieves higher robustness by not having to deal with many
AST synthesis scenarios.
Closes #8554.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
| |
Integer mul(matrix, matrix) and mul(vector, matrix) are not
disambiguated between __BuiltinIntegerType and __BuiltinLogicalType,
emitting an ambiguous call compilation error.
Use the OverloadRank attribute to prefer the IntegerType overload over
the LogicalType overload.
Fixes #8424
|
| |
|
|
|
|
|
| |
This is important for SPIR-V targets that need to know if a texture is
designated as a depth texture or not (for example WebGPU).
I didn't change the default behavior for when isShadow() is not set,
since I didn't want to make the change too invasive.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
- [x] Add `gl_DrawID` property declaration in
`source/slang/glsl.meta.slang` similar to `gl_BaseVertex` and
`gl_BaseInstance`
- [x] Ensure it maps to `SV_DrawIndex` semantic
- [x] Add extension requirements in `slang-ir-glsl-legalize.cpp` (GLSL
460 and GL_ARB_shader_draw_parameters)
- [x] Create test case to verify SPIRV output contains `DrawIndex`
decoration
- [x] Run formatting script before committing
- [x] Verify existing tests still pass
- [x] Remove accidentally committed temp files (1.glsl, 3.spv-asm)
- [x] Update .gitignore to prevent similar files from being committed
<!-- START COPILOT CODING AGENT SUFFIX -->
<details>
<summary>Original prompt</summary>
>
> ----
>
> *This section details on the original issue you should resolve*
>
> <issue_title>Missing direct usage support for gl_DrawID</issue_title>
> <issue_description># Problem Description
> Because DirectX does not have support for DrawIndex like vulkan does,
i add the support in d3d12 by emulating it and storing it in a root
constant. In vulkan though, i should be able to use it directly, but
because now i store it as a global variable, i cannot use the
SV_DrawIndex semantic either.
>
> # Preferred Solution
> I'd like to be able to use gl_DrawID like you how can use
gl_BaseVertex and gl_BaseInstance.
> </issue_description>
>
> <agent_instructions>You can just edit glsl.meta.slang and add a
declaration for gl_DrawID in a similar way to gl_BaseVertex, and map it
to SV_DrawIndex.
>
> Create a `//TEST:SIMPLE(filecheck=SPIRV): -target spirv` test case,
and use filecheck to verify the output contains the correct SPIRV
`DrawIndex` decoration.</agent_instructions>
>
> ## Comments on the Issue (you are @copilot in this section)
>
> <comments>
> </comments>
>
</details>
Fixes shader-slang/slang#8548
<!-- START COPILOT CODING AGENT TIPS -->
---
✨ Let Copilot coding agent [set things up for
you](https://github.com/shader-slang/slang/issues/new?title=✨+Set+up+Copilot+instructions&body=Configure%20instructions%20for%20this%20repository%20as%20documented%20in%20%5BBest%20practices%20for%20Copilot%20coding%20agent%20in%20your%20repository%5D%28https://gh.io/copilot-coding-agent-tips%29%2E%0A%0A%3COnboard%20this%20repo%3E&assignees=copilot)
— coding agent works faster and does higher quality work when set up for
your repo.
---------
Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Closes #7606.
When Slang compile for a bindful target, we will run the resource type
legalization pass to hoist resource typed struct fields outside of the
struct type and define them as global parameters and passing them around
via dedicated function parameters.
When we compile for a bindless target, we don't run this pass.
However, Metal is a hybrid bindful and bindless target. We need to run
type legalization for the constant buffer, but skip type legalization
for parameter block.
The previous attempt to support this behavior is to hack the type
legalization pass to return `LegalVal::simple` when it sees a
`ParameterBlock<T>`. However, whenever the code is accessing
`parameterBlock.someNestedField`, the type of the nested field may get a
`LegalType::tuple`, and now we will run into inconsistent scenarios
where we have a `LegalVal::simple` on the operand val, and but the
legalization logic is expecting that val to be a `LegalType::tuple`.
This breaks a lot of assumptions and invariants in the type legalization
pass, resulting unstable/fragile behavior.
To systematically solve this problem, this change generalizes the
existing legalize buffer element type pass to translate
`ParameterBlock<Texture2D>` (and similar cases) to
`ParameterBlock<Texture2D.Handle>`. So that such parameter block will
always be legalized to `LegalType:::simple` during type legalization,
and we will never run into any inconsistent cases. This allowed us to
get rid of the hacky logic in the type legalization pass to try to
workaround the inconsistencies.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Note that while this change touched a large numer of files, there are no
changes to functionality being made here. The only things being done are
renaming various symbols and, in a few cases, updating or adding
comments for consistency with the new names.
The core of the naming changes are:
* Most things named to refer to `OutType` (e.g., `IROutType`,
`IRBuilder::getOutType()`, etc.) have been consistently renamed to refer
to `OutParamType`, to emphasize that the relevant AST/IR node types are
only intended for use to represent `out` parameters.
* The same change as described above for `OutType` is also made for
`RefType`, which becomes `RefParamType` in most cases. One mess that
this exposes is the way that the `ExplicitRef<T>` type in the core
module currently lowers to `IRRefParamType`. This change sticks to the
rule of not making functional changes, so that mess is left as-is for
now.
* Names referring to `InOutType` have been changed to instead refer to
`BorrowInOutType`. The intention with this naming change is to emphasize
that the Slang rules for `inout` are semantically those of a borrow (or
at least our interpretation of what a borrow means).
* Names referring to `ConstRefType` have been changed to instead refer
to `BorrowInType`. This change starts work on clarifying that the
existing `__constref` modifier was never intended to be a read-only
analogue of `__ref`, and instead is the input-only analogue of `inout`.
* The `ParameterDirection` enum type has been changed to
`ParamPassingMode`, to reflect the fact that the concept of "direction"
fails to capture what is actually being encoded, particularly once we
have modes beyond simple `in`/`out`/`inout`.
While this change does not alter behavior in any case (the user-exposed
Slang language is unchanged), it is intended to set up subsequence
changes that will work to make the handling of these types in the
compiler more nuanced and correct. Breaking this part of the change out
separately is primarily motivated by a desire to minimize the effort for
reviewers.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
| |
Close #8568.
The root cause of this issue is that when the struct is indirectly
inherited from IDifferentiable type, we will not check the reference of
the DerivativeMember attribute. This PR fixes this issue by checking the
DerivativeMember attribute right before synthesize the requirement
methods of IDifferentiable interface.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
ByteAddressBuffer with DescriptorHandle (#8252)
- [x] Fix segmentation fault in wrapConstantBufferElement for
DescriptorHandle types
- [x] Split DescriptorKind.Buffer into ConstantBuffer and StorageBuffer
- [x] Update binding enums with descriptive names (ConstantBuffer_Read,
StorageBuffer_Read, etc.)
- [x] Update resource type mappings for correct binding assignments
- [x] Update template logic to handle ConstantBuffer and StorageBuffer
kinds separately
- [x] Update tests to reflect correct binding assignments
- [x] Split DescriptorKind.TexelBuffer into UniformTexelBuffer and
StorageTexelBuffer
- [x] Update TextureBuffer<T> to use UniformTexelBuffer kind
- [x] Update _Texture extension to determine texel buffer kind based on
access mode
- [x] Update test desc-handle-1.slang to handle new DescriptorKind enum
cases
<!-- START COPILOT CODING AGENT TIPS -->
---
✨ Let Copilot coding agent [set things up for
you](https://github.com/shader-slang/slang/issues/new?title=✨+Set+up+Copilot+instructions&body=Configure%20instructions%20for%20this%20repository%20as%20documented%20in%20%5BBest%20practices%20for%20Copilot%20coding%20agent%20in%20your%20repository%5D%28https://gh.io/copilot-coding-agent-tips%29%2E%0A%0A%3COnboard%20this%20repo%3E&assignees=copilot)
— coding agent works faster and does higher quality work when set up for
your repo.
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
|
|
| |
- Fix bug parsing multiple link-time structs on the same line. Closes
#8553.
- Fix bug parsing anonymous struct type as function return type in
modern syntax. Closes #8558
- Support semantics on modern style param/var declarations.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#8547)
This allows us to specialize functions whose argument is a sub element
of a constant buffer, instead of being only applicable to entire buffer
element. Closes #8421.
This change also implements a proper heuristic to determine when to
specialize the calls and defer the buffer loads.
This PR addresses a pathological case exposed in
`slangpy\slangpy\benchmarks\test_benchmark_tensor.py`, which used to
take 27ms to finish, and now takes 1.25ms.
For example, given:
```
struct Bottom
{
float bigArray[1024];
[mutating]
void setVal(int index, float value) { bigArray[index] = value; }
}
struct Root
{
Bottom top[2];
[mutating]
void setTopVal(int x, int y, float value)
{
top[x].setVal(y, value);
}
}
RWStructuredBuffer<Root> sb;
[shader("compute")]
[numthreads(1, 1, 1)]
void compute_main(uint3 tid: SV_DispatchThreadID)
{
sb[0].setTopVal(1, 2, 100.0f);
}
```
We are now able to specialize the call to `setTopVal` into:
```
void compute_main(uint3 tid: SV_DispatchThreadID)
{
setTopVal_specialized(0, 1, 2, 100.0f);
}
void setTopVal_specialized(int sbIdx, int x, int y, float value)
{
Bottom_setVal_specialized(sbIdx, x, y, value);
}
void Bottom_setVal_specialized(int sbIdx, int x, int y, float value)
{
sb[sbIdx].top[x].bigArray[y] = value;
}
```
And get rid of all unnecessary loads. Achieving this requires a
combination of function call specialization and buffer-load-defer pass.
The buffer-load-defer pass has been completely rewritten to be more
correct and avoid introducing redundant loads.
This PR also adds tests to make sure pointers, bindless handles, and
loads from structured buffer or constant buffers works as expected.
|
| |
|
|
|
|
|
|
|
| |
legalization pass. (#8567)
This is crash that be triggered by providing custom
`getDescriptorFromHandle` and use it to return access a
ByteAddressBuffer from a bindless handle.
Closes #8355.
|
| |
|
|
|
|
|
|
|
|
|
|
| |
Enables all tests/metal/ tests that can be easily enabled.
These tests were not originally designed as render tests; they are
generally being enabled for pipecleaning purposes, and will not be
rigorously testing the corresponding funcitonality.
Where they cannot be enabled as render tests, and a metallib test wasn't
already enabled, a metallib test was enabled instead (where possible).
Fixes #7892
|
| |
|
|
|
|
|
|
|
|
| |
Fixes #8439
When checked, generic type equality constraints types are now in a
canonical order, allowing for a commutative type equality operator.
---------
Co-authored-by: Mukund Keshava <mkeshava@nvidia.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
packing/unpacking. (#8526)
Part of the effort to improve the performance of generated SPIRV code.
The existing lower-buffer-element-type pass works by loading the entire
buffer element content from memory, and translate it to logical type
stored in a local variable at the earliest reference of a buffer handle.
This means that is can generate inefficient code that reads more than
necessary.
Consider this example:
```
struct BigStruct { bool values[1024]; }
ConstantBuffer<BigStruct> cb;
void test(BigStruct v)
{
if (v.values[0]) { printf("ok"); }
}
[numthreads(1,1,1)]
void computeMain()
{
test(cb);
}
```
In IR, the `computeMain` function before lower-buffer-element-type pass
is something like following:
```
func test:
%v = param : BigStruct
%barr = fieldExtract(%v, "values")
%element = elementExtract(%barr, 0)
... // uses %element
func computeMain:
%v = load(cb)
call %test %v
```
The existing lower-buffer-element-type pass will rewrite the bool array
in `BigStruct` into `int` array so it is legal in SPIRV. However, it
does so by inserting the translation on the first `load` of the constant
buffer:
```
struct BigStruct_std430 {
int values[1024];
}
var cb : ConstantBuffer<BigStruct_std430>;
func computeMain:
%tmpVar : var<BigStruct>
call %unpackStorage(%tmpVar, cb)
%v : BigStruct = load %tmpVar
call %test %v
```
This means that the entire array will be loaded and translated to int,
before calling `test`, which only uses one element. It turns out that
the downstream compiler isn't always able to optimize out this
inefficient translation/copy.
This PR completely rewrites the way buffer-element-type lowering is
handled to avoid producing this inefficient code. It works in two parts:
first we turn on the `transformParamsToConstRef` pass for SPIRV target
as well, so we will translate the `test` function to take the `v`
parameter as `constref`. The second part is a redesigned
buffer-element-type pass that defers the storage-type to logical-type
translation until a value is actually used by a `load` instruction.
In this example, after `transformParamsToConstRef`, the IR is:
```
func test:
%v = param : ConstRef<BigStruct>
%barr = fieldAddr(%v, "values")
%elementPtr = elementAddr(%barr, 0)
%element = load(%elementPtr)
... // uses %element
func computeMain:
call %test %cb
```
The new `buffer-element-type-lowering` pass will take this IR, and
insert translation at latest possible time across the entire call graph,
and translate the IR into:
```
func test:
%v = param : ConstRef<BigStruct_std430>
%barr = fieldAddr(%v, "values")
%elementPtr : ptr<int> = elementAddr(%barr, 0)
%element_int = load(%elementPtr)
%element = cast(%element_int) : %bool
... // uses %element
func computeMain:
call %test %cb
```
In this new IR, there is no longer a load and conversion of the entire
array.
See new comment in `slang-ir-lower-buffer-element-type.cpp` for more
details of how the pass works.
This PR also address many other issues surfaced by turning on
`transformParamsToConstRef` pass on SPIRV backend.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
| |
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.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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
|
| |
|
|
|
|
|
|
|
|
| |
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
|
| |
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
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.
|
