| Commit message (Collapse) | Author | Age |
| |
|
| |
Closes https://github.com/shader-slang/slang/issues/8378
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
…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>
|
| |
|
| |
The test can be enabled
|
| | |
|
| | |
|
| |
|
|
|
|
|
|
|
| |
Changes default for render-test to sm_6_5.
Since sm_6_5 is the new default, remove the -use-dxil option, add
-use-dxcb option
Remove -use-dxil option from all test cases.
Add -use-dxcb to two tests that needed it.
Fixes #7611
|
| | |
|
| | |
|
| | |
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
expressions in legacy mode (#7984)
This PR implements a warning system to help users identify potentially
unintended comma operator usage in expressions. The comma operator can
be confusing when used in contexts like variable initialization where
users might have intended to use braces for initialization instead.
## Problem
The following code compiles without error but is likely not written as
intended:
```slang
float4 vColor = (0.f, 0.f, 0.f, 1.f); // Uses comma operators, evaluates to 1.f
```
The intended code should use braces:
```slang
float4 vColor = {0.f, 0.f, 0.f, 1.f}; // Proper initialization
```
## Solution
Added a new warning diagnostic (`commaOperatorUsedInExpression`, ID:
41024) that warns when comma operators are used in expressions, with
exemptions for contexts where they are commonly intended:
- **For-loop side effects**: `for (int i = 0; i < 10; i++, x++)` - no
warning
- **Expand expressions**: `expand(f(), g(each param))` - no warning
- **Slang 2026+ mode**: `let m = (1,2,3)` creates tuples - no warning
- **All other expressions**: `float4 v = (a, b, c, d)` and `return a, b`
- warns for each comma
## Implementation Details
- Added context tracking in `SemanticsContext` with
`m_inForLoopSideEffect` flag
- Modified `visitForStmt` to use special context when checking side
effect expressions
- Added comma operator detection in `visitInvokeExpr` for `InfixExpr`
nodes
- Added language version check using `isSlang2026OrLater()` to disable
warnings in Slang 2026+ mode where parentheses create tuples
- Performance optimization: language version check is hoisted to avoid
unnecessary casting
- Warning can be suppressed using `-Wno-41024` command line flag
## Test Coverage
Added comprehensive test cases using filecheck format that verify:
- Warnings are generated for comma operators in variable initialization
(legacy mode only)
- Warnings are generated for comma operators in return statements
(legacy mode only)
- Warnings are generated for comma operators in general expressions
(legacy mode only)
- No warnings for comma operators in for-loop side effects
- No warnings in Slang 2026+ mode where parentheses create tuples
- Warning suppression works correctly
Example output (legacy mode):
```
warning 41024: comma operator used in expression (may be unintended)
float4 vColor = (0.f, 0.f, 0.f, 1.f);
^
warning 41024: comma operator used in expression (may be unintended)
return a *= 2, a + 1;
^
```
Fixes #6732.
<!-- START COPILOT CODING AGENT TIPS -->
---
💬 Share your feedback on Copilot coding agent for the chance to win a
$200 gift card! Click
[here](https://survey.alchemer.com/s3/8343779/Copilot-Coding-agent) to
start the survey.
---------
Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Co-authored-by: aidanfnv <198290069+aidanfnv@users.noreply.github.com>
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: aidanfnv <aidanf@nvidia.com>
Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com>
|
| | |
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Add emit cases for WGSL and GLSL
* Fix compilation warnings
Modify short cutting test to reflect change in emit logic
Lower matrix for metal as well
Add emit matrix logic for metal
Fix compiler warning
Brace initializer for lowered matrices
Fix compiler warnings
* Tests for metal
* Fix mult, any, and determinant
* Fix matrix-matrix multiplication
* Fix mat mul to be element-wise
* Fix compiler warning
* Move makeMatrix to legalization
* Move unary and binary arithmetic operator lowering to legalization
* Remove emit logic and move final comparison operators to legalization
* Handle vector/matrix negation for WGSL
* Restore older SPIR-V emit logic
* Address PR comments
* Revert to zero minus for negation
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(#7843)
* Fix 64-bit val lowering for metal
* Add ByteAddressBuffer load/store 64-bit tests
* Handle Store/Load ptr types
* Use bitcast for non-pointer typers
* format code (#7966)
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Add tests for expected behaviour
* Allow matrix types in logical or/and
* Legalize int/bool matrix types and construction with makeMatrix
* Legalize uint matrices and operations
* Limit testing to only SPIRV
* Better tests for int and bool
* Add test for uint
* Remove GLSL tests
* Remove old test for diagnosing int matrices
* Emit SPIRV directly in tests
* format code
* Address PR comments
* Improve testing
* Address PR comments
* format code
* Add tests for matrix intrinsic operations
* Move matrix lowering to dedicated legalization pass
* Fix compiler warning
* Remove signal again
* Reorder matrix and vector legalization
* Fix formatting
* Add shift and comparison tests
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Fix 1D texture reads in CUDA target
Fixes #7570: 1D surface writes don't work
The issue was that the Load function for read-only textures (hlsl.meta.slang lines 3629-3656)
only supported 2D and 3D textures for CUDA targets, causing 1D texture reads to fall through
to <invalid intrinsic>. This affected the srcTexture[tid.x] read operation in the reproduction case.
Changes:
- Updated static_assert to include SLANG_TEXTURE_1D support
- Added tex1DArrayfetch_int<T> for 1D array texture reads
- Added tex1Dfetch_int<T> for regular 1D texture reads
🤖 Generated with [Claude Code](https://claude.ai/code)
Co-Authored-By: Mukund Keshava <mkeshavaNV@users.noreply.github.com>
* Add 1D texture read support for CUDA target
- Add tex1Dfetch_int template specializations for float2, float4, uint, uint2, uint4
- Remove TODO comment about 1D PTX not being supported
- Enable 1D texture test in texture-subscript-cuda.slang
- Fix assembly code issues in original template specializations
🤖 Generated with [Claude Code](https://claude.ai/code)
Co-Authored-By: Mukund Keshava <mkeshavaNV@users.noreply.github.com>
* Update slang-cuda-prelude.h
* Fix texture3d ptx issue
* undo 1D texture changes
* Update hlsl.meta.slang
* Update hlsl.meta.slang
* Update hlsl.meta.slang
* Update hlsl.meta.slang
* Extend texture-subscript-cuda.slang test with uint and int format variants
Add test cases for newly supported texture formats in CUDA:
- 2D textures with uint, uint2, uint4
- 2D textures with int, int2, int4
- 3D textures with uint, uint2, uint4
- 3D textures with int, int2, int4
This ensures the texture subscript operations work correctly for all
the format variants added in the CUDA texture fixes.
Co-authored-by: Mukund Keshava <mkeshavaNV@users.noreply.github.com>
* update expected file
---------
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: Mukund Keshava <mkeshavaNV@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Fix IEEE 754 NaN comparisons in constant folding
Added proper NaN handling in SCCP optimization pass to follow IEEE 754 standard:
- NaN \!= any value returns true
- All other NaN comparisons return false
- Added double precision NaN detection support
- Fixed type detection to check operands instead of result type
* Avoid differentiating NaN and non-NaN cases
* format code (#76)
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
|
| |
|
|
|
|
|
|
|
|
|
| |
* fix #7554
* format code
* test ms and non ms texture
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
fixes: [#7143](https://github.com/shader-slang/slang/issues/7143)
fixes: [#7146](https://github.com/shader-slang/slang/issues/7146)
Goal of PR:
* This PR is part of the larger #7115 refactor to how dynamic dispatch works.
* The first step is to add the `-std <std-revision>` flag.
* The second step is to provide basic `dyn` keyword support in AST. This does not include `varDecl` support since most of these interactions require `some` keyword support.
Future PR(s) goal:
* Support `some` keyword in AST. With this we will also implement all varDecl interactions between `dyn` and `some`.
* Add IR support for `some` and `dyn`.
Breakdown of PR:
* most of the logic is in `validateDyn.*`. This was done so that in the future when we implement more features we will have an easy time removing/adding restrictions to `dyn` interfaces.
Breaking changes:
* As per spec (https://github.com/shader-slang/spec/pull/14/files), any type conforming to a `dyn` interface errors if member list contains one of the following: opaque type, non copyable type, or unsized type.
* Due to the breaking change, the test `tests\compute\dynamic-dispatch-bindless-texture.slang` is incorrect. This has been fixed.
|
| |
|
|
|
|
|
|
|
| |
* Fix HLSL ByteAddressBuffer Load* parameter integer type
* Fix tests
* Fix load with alignment function signature clash
* Fix LoadAligned tests
|
| |
|
|
|
|
|
|
|
|
|
| |
* Add checking for hlsl register semantic.
* Fix.
* Fix test.
* Fix switch error.
* Fix tests.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* WiP: Add more formats for texture reads
* fix test
* format code
* add float2/float4 versions for 1D and 3D as well
* fixed review comment
* fix review comments
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
| |
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* cuda: Add support for subscript operator
This CL adds support for the subscript operator for Read Only
textures in cuda. Also adds a test for this.
Fixes #6781
* format code
* fix review comments
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* update slang-rhi submodule
* slang-rhi API changes
* disable agility sdk
* fix texture creation
* update formats in tests
* Extent3D rename
* use 1 mip level for 1D textures for Metal
* fix texture upload
* update to latest slang-rhi
* update slang-rhi
* format code
* update slang-rhi
* do not run texture-intrinsics test on metal
* update slang-rhi
* deal with failing tests
* fix more tests
* update slang-rhi
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Simon Kallweit <simon.kallweit@gmail.com>
|
| |
|
|
|
|
|
|
| |
Fixes #6171
This commit adds logic for reporting double support
to the d3d11 backend, for running tests on GPUs that
do not support D3D11_FEATURE_DOUBLES, and add checks
for that support to tests that require the feature.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Add GetDimensions support for CUDA
This CL adds GetDimensions support for cuda by using the PTX
instructions. Currently, PTX only supports getting width, height and
depth.
This CL also adds a new test to test this support.
Fixes #5139
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
| |
- Add the diagnostic messages, and code to emit them
- Add some tests
This helps to address issue #6183.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Add a simple interface parameter test
Since there's no documentation, it's nice to have a simple test case in order to
experiment with this feature of the testing framework.
* Add shader entry point attributes to tests
* Fix specialization arguments for tests
- Add some missing arguments
- Rremove one extraneous argument.
* Stop using deprecated compile request in render-test
Use a session object instead of the deprecated compile request object.
This closes issue #4760.
|
| |
|
|
|
|
|
| |
* Update SPIRV-Tools and fix new validation errors.
* Implement pointers for glsl target.
* Reworked packStorage/unpackStorage code gen to operate on pointers rather than values.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
* Remove tests/compute/dump-repro
The -load-repro option is no longer maintained.
This helps to address issue #4760.
* Rename ShaderCompilerUtil::Output::session to globalSession
* Remove the load-repro codepath
* Lifetime bugfix
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* remove unused resource
* define buffer data
* add vs2022 build presets
* update slang-rhi API usage
* update slang-rhi
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
|
|
|
|
|
|
|
| |
* Use and() and or() functions for logical-AND and OR
With this commit, Slang will emit function calls to `and()` and `or()`
for the logical-AND and logical-OR when the operands are non-scalar and
the target profile is SM6.0 and above. This is required change from
SM6.0.
For WGSL, there is no operator overloadings of `&&` and `||` when the
operands are non-scalar. Unlike HLSL, WGSL also don't have `and()` nor
`or()`. Alternatively, we can use `select()`.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* SP004: implement initialize list translation to ctor
- We synthesize a member-wise constructor for each struct follow
the rules described in SP004.
- Add logic to translate the initialize list to constructor invoke
- Add cuda-host decoration for the synthesized constructor
- Remove the default constructor when we have a valid member init constructor
- Disable -zero-initialize option, will re-implement it in followup (#6109).
- Fix the overload lookup issue
When creating invoke expression for ctor, we need to call
ResolveInvoke() to find us the best candidates, however
the existing lookup logic could find us the base constructor
for child struct, we should eliminate this case by providing
the LookupOptions::IgnoreInheritance to lookup, this requires
us to create a subcontext on SemanticsVisitor to indicate that
we only want to use this option on looking the constructor.
- Do not implicit initialize a struct that doesn't have explicit default
constructor.
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
* Implement anyvalue marshalling for 8-bit integers
* Fix missing offset from int8/uint8 case
* Disable anyvalue 8-bit test for DXIL
Because it doesn't support 8-bit values anyway.
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Initial implementation of `ResourcePtr<T>`.
* Update docs
* Fix build error.
* Add more discussion.
* Update documentation.
* Update TOC.
* Fix.
* Fix.
* Add test case for custom `getResourceFromBindlessHandle`.
* Add namehint to generated descriptor heap param.
* Fix.
* Fix.
* format code
* Rename to `DescriptorHandle`, and add `T.Handle` alias.
* Fix compiler error.
* Fix.
* Fix build.
* Renames.
* Fix documentation.
* Documentation fix.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
| |
Close #5470
* implement bitcast for 64-bit date type
* Move 'ensurePrelude' to base class to remove duplication
* Assert on 'double' type for Metal target, as Metal doesn't have 'double' support
|
| | |
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Respect explicit bindings in wgsl emit.
* Implement explciit binding generation for metal and wgsl.
* Update toc.
* Fix warnings in tests.
* Fix tests.
---------
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Add buffer swizzle store test using structured buffers
This helps to address #5612
* WGPU: Add explanation for tests/bugs/buffer-swizzle-store.slang
This closes #5612.
* Manual formatting fix
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Closes https://github.com/shader-slang/slang/issues/5067
New tests, covering what's declared supported in the WGSL support docs
- tests/wgsl/semantic-coverage.slang
- tests/wgsl/semantic-depth.slang
- tests/wgsl/semantic-dispatch-thread-id.slang
- tests/wgsl/semantic-group-id.slang
- tests/wgsl/semantic-group-index.slang
- tests/wgsl/semantic-group-thread-id.slang
- tests/wgsl/semantic-instance-id.slang
- tests/wgsl/semantic-is-front-face.slang
- tests/wgsl/semantic-position.slang
- tests/wgsl/semantic-sample-index.slang
- tests/wgsl/semantic-vertex-id.slang
WGSL enabled existing tests:
- tests/compute/compile-time-loop.slang
- tests/compute/constexpr.slang
- tests/compute/discard-stmt.slang
- tests/metal/nested-struct-fragment-input.slang
- tests/metal/nested-struct-fragment-output.slang
- tests/metal/nested-struct-multi-entry-point-vertex.slang
- tests/metal/no-struct-vertex-output.slang
- tests/metal/sv_target-complex-1.slang
- tests/metal/sv_target-complex-2.slang
- tests/bugs/texture2d-gather.hlsl
- tests/render/cross-compile-entry-point.slang
- tests/render/nointerpolation.hlsl
- tests/render/render0.hlsl
- tests/render/cross-compile0.hlsl
- tests/render/imported-parameters.hlsl
- tests/render/unused-discard.hlsl
Can't be enabled due to missing wgsl features
- tests/compute/texture-sampling.slang
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
| |
* Don't treat StrcturedBuffer<IFoo> as a specializable param.
* Fix RHI.
|
| | |
|
| |
|
|
|
|
|
|
|
| |
Some tests are now passing and are enabled.
Other tests are still failing, but are given comments categorizing the failures.
Tests in the 'Not supported in WGSL' category are also removed from the expected failures
list. (Though they are still kept disabled for WebGPU, of course.)
This closes #5519.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Limit number of MIP levels on 1d textures to be 1
This avoids running into a WebGPU limitation, and helps to address issue #4943.
* Update Slang-RHI to get WGPU fixes
* Enable WGPU texture sampling test
This helps to address issue #4943.
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Add new texture sampling test for WebGPU
There are no 1d array textures in WGSL, so
add texture-sampling-no-1d-arrays.slang based on texture-sampling.slang, but without
1d texture arrays.
This helps to address issue #4943.
* Insert needed conversion when querying texture attributes in WGSL
This helps to address issue #4943.
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
* Update Slang-RHI to get WGPU backend fixes
* render-test: Use device local memory type for vertex buffers
This helps to avoid https://github.com/shader-slang/slang-rhi/issues/104
* Fix bug in WGSL emitter layout code.
There was a "kinds" vs. "kind flags" mismatch, and also getBindingOffsetForKinds was not
being used.
This patch enables a bunch of tests for WGPU.
This helps to address issue #4943.
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
| |
|
| |
This closes issue #5505.
|