summaryrefslogtreecommitdiffstats
path: root/tests/compute
Commit message (Collapse)AuthorAge
* implement dot products for 1 vectors (#8599)Ellie Hermaszewska2025-10-10
| | | Closes https://github.com/shader-slang/slang/issues/8378
* Fix DerivativeGroupQuadsKHR workgroup size validation for texture sampling ↵Lujin Wang2025-10-08
| | | | | | | | | | | | | | | | | | | | | | | (#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>
* Rewriting the lower-buffer-element-type pass to avoid unnecessary ↵Yong He2025-09-30
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 (#8340)Harsh Aggarwal (NVIDIA)2025-09-18
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Fix CUDA global variable initialization with constructor calls Resolves CUDA compilation failure where global variables with struct constructor initialization generated illegal `__device__` variable runtime initialization. **Problem:** ```cuda // Generated invalid CUDA code: __device__ static const Stuff_0 gStuff_0 = Stuff_x24init_0(args...); // Error: "dynamic initialization is not supported for a __device__ variable" Root Cause Discovered: Through extensive debugging, found that moveGlobalVarInitializationToEntryPoints pass only handled kIROp_GlobalVar instructions, but global constants with constructor calls appeared as kIROp_Call instructions at module scope. Solution: 1. IR Pipeline Fix: Extended moveGlobalVarInitializationToEntryPoints to detect and transform module-level constructor calls into proper global variables with entry-point initialization 2. Field Access Fix: Enhanced kIROp_FieldExtract logic to emit correct -> syntax for pointer types and address-of operations 3. Constructor Emission: Added CUDA-specific handling for constructor calls Architecture: - Transforms let %gStuff = call %Constructor(...) into kernel context initialization - Moves runtime initialization from global scope to entry-point execution - Follows CUDA best practices for global state management Files: - source/slang/slang-ir-explicit-global-init.cpp: Extended IR transformation pass - source/slang/slang-emit-c-like.cpp: Enhanced field access and foldable value logic - source/slang/slang-emit-cuda.cpp: Added CUDA-specific field extraction handling Result: // Now generates proper CUDA code: struct KernelContext_0 { Stuff_0 gStuff_1; }; // Runtime initialization in entry point: kernelContext_1.gStuff_1 = constructor_call(); Fixes: tests/compute/type-legalize-global-with-init.slang
* CUDA: Fix compiler crash with unsized array field - nonuniformres-as-… (#8380)Harsh Aggarwal (NVIDIA)2025-09-10
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | …function-parameter.slang #8315 Root Cause: CUDA compilation crashed with `assert failure: !seenFinalUnsizedArrayField` because unsized arrays like `RWStructuredBuffer<uint> globalBuffer[]` were not the final field in generated parameter structs, violating the layout constraint in slang-ir-layout.cpp. Fix: Extended `collectGlobalUniformParameters` to automatically reorder struct fields for CUDA targets - regular fields first, unsized arrays last. Other targets preserve original order. Impact: - Enables CUDA support for nonuniform resource indexing as function parameters - Zero impact on existing GLSL/HLSL/SPIRV targets - Automatic handling - no manual parameter reordering required Files: slang-emit.cpp, slang-ir-collect-global-uniforms.cpp/.h, test file --------- Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com> Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
* Fix #8314 - Enable tests/compute/texture-subscript.slang for CUDA (#8408)Harsh Aggarwal (NVIDIA)2025-09-09
| | | The test can be enabled
* Fix#8085: Batch-9: Enable cuda tests (#8269)Harsh Aggarwal (NVIDIA)2025-09-03
|
* Fix#8086: Batch-10: Enable cuda tests (#8270)Harsh Aggarwal (NVIDIA)2025-09-03
|
* render-test: Change D3D12 default to sm_6_5 (#8320)James Helferty (NVIDIA)2025-09-02
| | | | | | | | | 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
* Fix#8084: Batch-8: Enable cuda tests (#8268)Harsh Aggarwal (NVIDIA)2025-08-25
|
* Fix#8083: Batch-7: Enable cuda tests (#8267)Harsh Aggarwal (NVIDIA)2025-08-25
|
* Fix#8082: Batch-6: Enable cuda tests (#8266)Harsh Aggarwal (NVIDIA)2025-08-25
|
* Add warning for comma operators used outside for-loops and expand ↵Copilot2025-08-07
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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>
* Enable compute/ dir which passes (#7991)Harsh Aggarwal (NVIDIA)2025-08-05
|
* Lowering unsupported matrix types for GLSL/WGSL/Metal targets (#7936)venkataram-nv2025-07-30
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * 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>
* Fix Metal invalid as_type cast for 64-bit RWByteAddressBuffer.Store values ↵Gangzheng Tong2025-07-29
| | | | | | | | | | | | | | | | | | | | (#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>
* Lower int/uint/bool matrices to arrays for SPIRV (#7687)venkataram-nv2025-07-18
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * 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 CUDA issues with texture reads and surface writes (#7780)Mukund Keshava2025-07-16
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * 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 (#7721)Jay Kwak2025-07-11
| | | | | | | | | | | | | | | | | * 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>
* Emit sample index when constructing a `OpImageTexelPointer` (#7563)ArielG-NV2025-06-30
| | | | | | | | | | | * fix #7554 * format code * test ms and non ms texture --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Initial `dyn` keyword support & `-lang 2026` compiler option (#7172)ArielG-NV2025-05-22
| | | | | | | | | | | | | | | | | | | | 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 (#7117)Darren Wihandi2025-05-16
| | | | | | | | | * 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. (#7118)Yong He2025-05-15
| | | | | | | | | | | * Add checking for hlsl register semantic. * Fix. * Fix test. * Fix switch error. * Fix tests.
* cuda: Add more formats for texture read/write (#7012)Mukund Keshava2025-05-12
| | | | | | | | | | | | | | | | | | * 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>
* Add a new capability hlsl_2018 that avoid using select/and/or (#7003)Jay Kwak2025-05-05
| | | Co-authored-by: Yong He <yonghe@outlook.com>
* Add subscript operator support in cuda (#6830)Mukund Keshava2025-04-30
| | | | | | | | | | | | | | | | | | | * 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 (#6587)Simon Kallweit2025-04-24
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * 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>
* Add slang-test check for D3D11 double support (#6761)aidanfnv2025-04-12
| | | | | | | | 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 (#6718)Mukund Keshava2025-04-01
| | | | | | | | | | | | | | | | | * 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 error diagnostic vectors and matrices with disallowed element types (#6596)Anders Leino2025-03-17
| | | | | | - Add the diagnostic messages, and code to emit them - Add some tests This helps to address issue #6183.
* Migrate render-test away from deprecated compile request API (#6514)Anders Leino2025-03-12
| | | | | | | | | | | | | | | | | | * 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. (#6511)Yong He2025-03-06
| | | | | | | * 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.
* Prepare for render test api migration (#6498)Anders Leino2025-02-28
| | | | | | | | | | | | | * 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
* update slang-rhi (shader object refactor) (#6251)Simon Kallweit2025-02-27
| | | | | | | | | | | | | | | * 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 (#6310)Jay Kwak2025-02-07
| | | | | | | | | | | | * 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()`.
* Feature/initialize list side branch (#6058)kaizhangNV2025-02-05
| | | | | | | | | | | | | | | | | | | | | | * 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 (#6059)Julius Ikkala2025-01-15
| | | | | | | | | | | | | * 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 SP#015 `DescriptorHandle<T>`. (#6028)Yong He2025-01-10
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * 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>
* Implement bitcast for 64-bit date type (#5895)kaizhangNV2024-12-17
| | | | | | | 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
* Support specialization constant on WGSL and Metal. (#5780)Yong He2024-12-06
|
* Implement explciit binding for metal and wgsl. (#5778)Yong He2024-12-06
| | | | | | | | | | | | | | | * 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>
* WGPU: Add new test and explanation for disabled test (#5731)Anders Leino2024-12-04
| | | | | | | | | | | | | | | | | * 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>
* Varying inputs and outputs for wgsl (#5669)Ellie Hermaszewska2024-12-02
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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. (#5645)Yong He2024-11-22
| | | | | * Don't treat StrcturedBuffer<IFoo> as a specializable param. * Fix RHI.
* Add explanation for disabled WGPU tests (#5628)Anders Leino2024-11-22
|
* Refresh of disabled WGPU tests (#5614)Anders Leino2024-11-21
| | | | | | | | | 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.
* Enable WGPU texture sampling test (#5617)Anders Leino2024-11-20
| | | | | | | | | | | | | | | * 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>
* Insert some casts for WGSL texture attribute queries (#5560)Anders Leino2024-11-14
| | | | | | | | | | | | | | | | | * 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>
* Various fixes to enable some WGSL graphics tests (#5548)Anders Leino2024-11-13
| | | | | | | | | | | | | | | | | | | | | * 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>
* Enable a bunch of WGPU tests (#5513)Anders Leino2024-11-07
| | | This closes issue #5505.