summaryrefslogtreecommitdiffstats
path: root/tests/hlsl-intrinsic
Commit message (Collapse)AuthorAge
* 8503 wgsl depth texture (#8645)Sami Kiminki (NVIDIA)2025-10-10
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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
* Disable branching subgroup test for WGSL (#8614)Jay Kwak2025-10-07
| | | | | | | | 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
* Prefer IntegerType over LogicalType integer matrix mul() overloads (#8426)pdeayton-nv2025-10-06
| | | | | | | | | | 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
* Enable metal tests (#8446)James Helferty (NVIDIA)2025-09-30
| | | | | | | | | | | | 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
* Prepare VulkanSDK release Oct 2025 (#8525)Jay Kwak2025-09-25
| | | | Related to - https://github.com/shader-slang/slang/issues/8519
* Disable a few WGSL testing using Wave functions in branches (#8386)Jay Kwak2025-09-05
| | | | | | | | | | WGSL requires Wave functions to be used only in uniform control flow. The latest compiler, Tint, started to error out when Wave functions are used in a dynamic control flow. This commit disables some of tests using Wave functions in dynamic branches. If possible, they are altered to call Wave functions in uniform control flows.
* Enable CUDA support for additional HLSL intrinsic tests (#8293)Harsh Aggarwal (NVIDIA)2025-09-04
| | | | | | | | | | | | | | | | | | | | | | | | Enable CUDA support for additional HLSL intrinsic tests by implementing missing functionality and fixing compiler bugs affecting CUDA targets. - Fix critical bug in InterlockedCompareStore64 where division used /4 instead of /8 for 64-bit types, causing incorrect memory addressing for all signed int 64_t atomics - Add signed int64_t atomic wrappers (atomicExch, atomicCAS) to CUDA prelu de that properly cast to/from unsigned types as required by CUDA's atomic API - Enable tests: atomic-intrinsics-64bit.slang - Implement CUDA support for QuadAny and QuadAll operations using warp shu ffle primitives (__shfl_sync with quad-level lane masking) - Add CUDA to quad_control capability definition in slang-capabilities.capdef - Add _slang_quadAny/_slang_quadAll helper functions to CUDA prelude - Enable tests: quad-control-comp-functionality.slang, subgroup-quad.slang --------- Co-authored-by: szihs <675653+szihs@users.noreply.github.com>
* 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 Metal 8-bit vector type names: emit char/uchar instead of int8_t/uint8_t ↵Copilot2025-08-26
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (#8223) The Metal backend was generating incorrect type names for 8-bit vector types, causing compilation failures when targeting Metal. According to the Metal specification, 8-bit vector types should be named `charN` and `ucharN` (e.g., `char2`, `uchar3`) rather than `int8_tN` and `uint8_tN`. ## Problem When compiling Slang code with 8-bit vector types for Metal, the compiler would emit: ```metal uint8_t2 _S8 = uint8_t2(uint8_t(0U), uint8_t(16U)); int8_t3 _S9 = int8_t3(int8_t(0), int8_t(16), int8_t(48)); ``` But the Metal compiler expects: ```metal uchar2 _S8 = uchar2(uint8_t(0U), uint8_t(16U)); char3 _S9 = char3(int8_t(0), int8_t(16), int8_t(48)); ``` This caused errors like: ``` error: unknown type name 'uint8_t2'; did you mean 'uint8_t'? ``` ## Solution Modified `MetalSourceEmitter::emitSimpleTypeImpl()` to emit the correct Metal-specific type names for 8-bit types: - `kIROp_Int8Type` now emits `char` instead of `int8_t` - `kIROp_UInt8Type` now emits `uchar` instead of `uint8_t` This change only affects the Metal backend and ensures that vector types like `int8_t2`, `uint8_t3`, etc. are correctly emitted as `char2`, `uchar3`, etc. ## Testing - Added a new test case `tests/metal/8bit-vector-types.slang` to verify the fix - Re-enabled the previously disabled Metal test in `tests/hlsl-intrinsic/countbits8.slang` - Updated `tests/metal/byte-address-buffer.slang` to expect the correct type names - Verified that existing Metal tests continue to pass Fixes #8211. <!-- START COPILOT CODING AGENT TIPS --> --- 💡 You can make Copilot smarter by setting up custom instructions, customizing its development environment and configuring Model Context Protocol (MCP) servers. Learn more [Copilot coding agent tips](https://gh.io/copilot-coding-agent-tips) in the docs. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: bmillsNV <163073245+bmillsNV@users.noreply.github.com>
* Fix#8082: Batch-6: Enable cuda tests (#8266)Harsh Aggarwal (NVIDIA)2025-08-25
|
* Fix#8081: Batch-5: Enable cuda tests (#8263)Harsh Aggarwal (NVIDIA)2025-08-25
|
* Fix#8080: Batch-4: Enable cuda tests (#8261)Harsh Aggarwal (NVIDIA)2025-08-25
|
* Add Metal support for WaveGetActiveMask and WaveActiveCountBits (#8218)Tianyu Li2025-08-20
| | | | | | | | | | | | | | | | | | | | ## Summary - Add Metal platform support for `WaveGetActiveMask()` and `WaveActiveCountBits()` wave intrinsics - Update capability requirements to include Metal platform for subgroup ballot operations - Implement Metal-specific intrinsic assembly using `simd_ballot()` and `simd_vote` APIs ## Changes - **source/slang/hlsl.meta.slang**: - Add Metal target case for `WaveGetActiveMask()` using `simd_ballot(true)` - Update capability requirements from `cuda_glsl_hlsl_spirv` to `cuda_glsl_hlsl_metal_spirv` for wave ballot functions - **source/slang/slang-capabilities.capdef**: - Add `metal` to `subgroup_ballot_activemask` capability alias
* Updated support to enable batch3 (#8219)Harsh Aggarwal (NVIDIA)2025-08-20
| | | | | | | | | Enable CUDA support for batch 3 tests - Enhanced wave operations with exclusive support - Added proper identity values for min/max operations - Fixed intrinsic name mapping issues - Updated test configurations Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
* Use 64bit int instead of emulation on metal (#8180)James Helferty (NVIDIA)2025-08-15
| | | | | | | | | | | | Metal's popcount prototype is `T popcount(T x)` but we want to use it to implement `countbits` where the prototype always returns `uint`. Using `popcount` directly would implicitly cast successfully to the 32-bit return value in all cases except when the argument is a 64-bit type. Thus, this change always explicitly casts the result to `$TR`, which should be one of the `uint[N]` types, and should always be able to hold the number of bits in the type. Addresses #6877
* Fix atomics error diagnostics (#8117)venkataram-nv2025-08-09
| | | | | | | Fixes #8116 --------- Co-authored-by: Jay Kwak <82421531+jkwak-work@users.noreply.github.com>
* Error if super-type capabilities are a super-set of sub-type (#7452)ArielG-NV2025-08-08
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Fixes: #7410 Changes: 1. super-type capabilities must be a super-set of sub-type capabilities (and support the same shader stages/targets) * InheritanceDecl visits super-type to inherit it's capabilities; validate InheritanceDecl capabilities against sub-type * visit all container decl's with a default case * clean up functionDeclBase visitor * Simplify `diagnoseUndeclaredCapability` by moving logic into capability checking (more correct*) 3. added changed behavior to documentation 4. fixed some incorrect capabilities 5. **we do not** diagnose capability errors on interface requirement-to-implementation if both lack explicit capability requirements. This change is to work around a slangpy regression (test case for the failing situation is in `tests\language-feature\capability\capability-interface-extension-1.slang`), Note: maybe for slang-2026 we don't do this? 6. requirement & implementation must support the same shader stage/target. This was changed because otherwise we can have cases where `X` inherits from `Y`, but `Y` is only expected to be used in `glsl` whilst `X` is expected to be used in `hlsl | glsl` 7. removed `tests/language-feature/capability/capabilitySimplification3.slang` because it tests nothing special (redundant) Note: not using rebase due to separate branches depending on this PR --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* 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>
* Add matrix select intrinsic (#7566)venkataram-nv2025-07-31
| | | | | | | | | | | | | | | | | | | | | | | * Add matrix select intrinsic * Fix hlsl test * Restrict matrix select to HLSL * Better test for HLSL side * Select route for GLSL/SPIRV * Exclude matrices from select legalization * Exclude CUDA from select test * Inline and move * format code --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Fix CUDA backend missing U32_firstbitlow implementation (#7921)Copilot2025-07-29
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * Initial plan * Add U32_firstbitlow implementation for CUDA and CPP backends Co-authored-by: bmillsNV <163073245+bmillsNV@users.noreply.github.com> * Add I32_firstbitlow and comprehensive testing for signed/unsigned firstbitlow Co-authored-by: bmillsNV <163073245+bmillsNV@users.noreply.github.com> * Convert firstbitlow test to use inline filecheck syntax Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com> * Add U32_firstbithigh and I32_firstbithigh implementations for CUDA and CPP backends Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com> * Update prelude/slang-cpp-scalar-intrinsics.h * Update prelude/slang-cpp-scalar-intrinsics.h * Update prelude/slang-cpp-scalar-intrinsics.h * Refactor Metal bit intrinsics to handle zero case correctly Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com> * Update slang-cuda-prelude.h remove fake links * Update hlsl.meta.slang * if -1, return -1 due to implicit hlsl rule * -1 or 0 is ~0u as per hlsl implictly * 0 or -1 as per hlsl * fix the math to map to hlsl * fix compile error * forgot `31 - clz` * format code (#7943) Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com> * Update source/slang/hlsl.meta.slang * Update source/slang/hlsl.meta.slang * Update source/slang/hlsl.meta.slang * Update source/slang/hlsl.meta.slang --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: bmillsNV <163073245+bmillsNV@users.noreply.github.com> Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com> Co-authored-by: csyonghe <2652293+csyonghe@users.noreply.github.com> Co-authored-by: ArielG-NV <aglasroth@nvidia.com> Co-authored-by: slangbot <ellieh+slangbot@nvidia.com> Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Enable tests for CUDA (#7593)Mukund Keshava2025-07-03
| | | | | | | | Enable intrinsic tests for cuda. Most of these tests were either disabled or just not enabled for cuda. Fixes #7592 Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
* Fix tuple AST & IR layout size queries (#7502)Julius Ikkala2025-06-26
| | | | | * Fix tuple AST & IR layout size queries * Don't peephole sizeof if size is still indeterminate
* Fix additional VVL violations (#7377)Gangzheng Tong2025-06-18
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * fix: add sampleCount and mipMaps to st2DMS_f32v4 Fix VUID-VkImageCreateInfo-samples-02257: The Vulkan spec states: If an OpTypeImage has an MS operand 1, its bound image must not have been created with VkImageCreateInfo::samples as VK_SAMPLE_COUNT_1_BIT * Fix VUID-VkShaderModuleCreateInfo-pCode-08740 Rename VK_KHR_COMPUTE_SHADER_DERIVATIVES_EXTENSION_NAME to VK_NV_COMPUTE_SHADER_DERIVATIVES_EXTENSION_NAME * fix: add sampleCount and mipMaps to st2DMS_f32v4 Fix VUID-VkImageCreateInfo-samples-02257: The Vulkan spec states: If an OpTypeImage has an MS operand 1, its bound image must not have been created with VkImageCreateInfo::samples as VK_SAMPLE_COUNT_1_BIT * Fix VUID-VkShaderModuleCreateInfo-pCode-08740 Rename VK_KHR_COMPUTE_SHADER_DERIVATIVES_EXTENSION_NAME to VK_NV_COMPUTE_SHADER_DERIVATIVES_EXTENSION_NAME * Fix VUID-vkCmdDispatch-None-06479 Use correct format for combined depth texture. * Fix VUID-vkCmdDispatch-format-07753 by setting format Parse filtering mode for sampler because the RGBA8* formats do not support linear filtering * Create MS texture type for sample count > 1 * Use different texture formats for depth compare and gather ops * Use clearTexture for init the data for MS textures
* Diagnose on use of struct inheritance. (#7419)Yong He2025-06-12
| | | | | | | | | | | | | * Diagnose on use of struct inheritance. * fix test. * Fix tests. * fix. --------- Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
* Address issues with GLSL style global in/out vars (#6669) (#6998)sricker-nvidia2025-06-06
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * Address issues with GLSL style global in/out vars (#6669) Asserts and segfaults were observed trying to compile a simple vertex shader like: ```` in int2 inPos; [shader("vertex")] main(uniform int2 test1, int2 test2, out float4 pos: SV_Position) void main() { // Bogus use of all input vars to prevent optimizing out. pos = float4(inPos.x, test1.x, test2.y, 0); } ```` Further investigation found that while replacing "uniform int2 test1" with "int2 test1" allowed for successful compilation, the resulting output shader would have overlapping location qualifiers. For example, compiling the above with "int2 test1" to glsl might give: ```` ... layout(location = 0) in ivec2 test1_0; layout(location = 1) in ivec2 test2_0; layout(location = 0) in ivec2 translatedGlobalParams_inPos_0; ... ```` This was because Slang does not actually support mixing GLSL style global in/out vars and entry point params. However, this is never checked for or noted in documentation. Slang source also assumes input shaders do not mix these and these assumptions ultimately led to the observed asserts and seg faults when using uniform entry point params. This change makes updates to throw an error when the compiler detects that it is trying to translate global in/out variables into entry point params when an entry point already contains parameters, allowing for compilation to fail gracefully. Certain tests have been updated to avoid mixing GLSL style global in/out vars and entry point params. This was mostly for tests that were using functions like WaveGetLaneIndex which use global in vars for certain platforms (see __builtinWaveLaneIndex). * Address issues with GLSL style global in/out vars - updates 1 (#6669) Update addresses review feedback to support mixing GLSL-flavored global in/out vars and entrypoint parameters when either all global in/out vars or all entry point params have a system value binding semantic. * Address issues with GLSL style global in/out vars - updates 2 (#6669) This update attempts to actually allow mixing GLSL style global in vars and entry point vars. Change attempts to recalculate offsets when adding the global input vars into the recreated entry point params layout. Additional updates were made to: -resolve further issues uncovered with entry point uniform params. -Address improper use of SV_DispatchThreadID in wave-get-lane-index.slang for metal. "thread_position_in_grid" is not supported for signed integer scalars or vectors. -Fix a spirv casting conflict due to the implementation of gl_PrimitiveID.get conflicting with PrimitiveIndex(). -Add a call to remove a global var in replaceUsesOfGlobalVar(). The global var is already replaced in this function and keeping it around can prevent it from being cleaned up by DCE if it still has decorations. * format code --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Add LSS intrinsics (#7200)Mukund Keshava2025-05-27
| | | | | | | | | | | | | * WiP: LSS intrinsics: initial commit * format code * Fix CI failures * Address review comment --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Implement shader execution reordering support for OptiX (#7211)Harsh Aggarwal (NVIDIA)2025-05-26
| | | | | | | | | | | | | | | | * Implement shader execution reordering support for OptiX Added OptiX backend support for Shader Execution Reordering (SER) features as outlined in issue #6647. This implementation: 1. Added CUDA target support for HitObject API 2. Implemented core SER functionality (TraceRay, MakeHit/Miss, Invoke) 3. Added OptiX-specific hit object handling functions 4. Added test case for OptiX SER functionality * format code --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Add full support for SPV_NV_shader_subgroup_partitioned (#7103)Darren Wihandi2025-05-25
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * Properly implement WaveMask* variants of WaveMultiPrefix* intrinsics * More partitioned intrinsics * More partitioned intrinsics and cleaned up non-prefixed WaveMask* implementations * Refactor HLSL WaveMultiPrefix* implementations * fix cap atoms * Clean up implementation * Add GLSL intrinsics and cleanup * Add tests * Fix affected capability test * Update and fix tests * Move expected.txt file * Refactor WaveMask* to call WaveMulti* * Refactor SPIRV/GLSL preamble code * Enable emit-via-glsl tests * remove wave_multi_prefix capability in favor of subgroup_partitioned * Update docs * Update cap atoms doc
* Make sizeof(T) & alignof(T) of generic types work as compile-time constants ↵Julius Ikkala2025-05-22
| | | | | | | | | | | (#7213) * Make sizeof(generic) work as compile-time constant * format code --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Fix broken -emit-spirv-via-glsl test option (#7091)sricker-nvidia2025-05-16
| | | | | | | | | | | | | | | | | | | | | | | | Fixes issue #6898 The -emit-spirv-via-glsl slang-test option has been broken for some amount of time. Tests that were using it were operating as if using -emit-spirv-directly, leading to many duplicated tests. After fixing the test option, there were an number of errors that appeared as a result. This change fixes the broken test option and the resulting test errors. Some of the test errors revealed some legitimate issues, such as: -The GLSL bitCount instrinsic only supports 32-bit integers and requires emulation for other bit widths. -Emitting GLSL 8-bit and 16-bit glsl integer types did not emit the proper extension requirements -Emitting GLSL and casting for 16-bit integers was missing a closing parenthesis. -Missing profile for GL_EXT_shader_explicit_arithmetic_types -Missing toType cases for UInt8/Int8 for the kIROp_BitCast case in tryEmitInstExprImpl.
* Remove readonly keyword from buffer pointer definitions (#7068)aidanfnv2025-05-14
| | | | | | For https://github.com/shader-slang/slang/issues/6880 This change removes the readonly keyword from buffer pointer definitions from the GLSL source emitter, to allow for mutable buffer pointers. Support for readonly will be readded when we add const pointer support later.
* cluster acceleration structure optix 6431 (#7028)Harsh Aggarwal (NVIDIA)2025-05-12
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * Add cluster geometry intrinsics for ray tracing - Added GetClusterID() method to HitObject class - Added CandidateClusterID() and CommittedClusterID() methods to RayQuery class - Added SPV_NV_cluster_acceleration_structure extension support - Added GL_NV_cluster_acceleration_structure extension support - Added test files for RayQuery and HitObject cluster methods Fixes #6431 * OpRayQueryGetIntersectionClusterIdNV - unrecognized spirv Disabling spirv backend for SPV_NV_cluster_acceleration_structure hlsl.meta.slang(18674): error 29100: unrecognized spirv opcode: OpRayQueryGetIntersectionClusterIdNV result:$$int = OpRayQueryGetIntersectionClusterIdNV &this $iCandidateOrCommitted; ^~~~~~ hlsl.meta.slang(18670): error 30019: expected an expression of type 'int', got 'void' return spirv_asm ^~~~~~~~~ ninja: build stopped: subcommand failed. * 6431 - Fix spirv opcode * Remove tests * Add relevant tests * Review - Simplify tests
* Add countbits 16-bit and 8-bit support (#6433) (#6897)sricker-nvidia2025-05-05
| | | | | | | | Change adds 16-bit and 8-bit support for countbits intrinsic. In cases where a backend's native counbits lacks support, support is emulated. New tests are added for 16-bit and 8-bit support. Additional testing added for 32-bit and minor updates made to 64-bit countbits.
* Disable an always failing test, partial resident test with dx12 (#6983)Jay Kwak2025-05-02
|
* Add fwidth_coarse and fwidth_fine functions (#6941)pdeayton-nv2025-05-01
| | | | | | | | Fixes #6940. Add new Slang fwidth_coarse and fwidth_fine functions, similar to GLSL's fwidthCoarse and fwidthFine. Move the implementation of the GLSL functions from glsl.meta.slang to hlsl.meta.slang. Update the existing spirv/fwidth.slang test with the new functions, and add a new hlsl-intrinsic/fragment-derivative.slang test to test HLSL, SPIR-V, and GLSL targets for the new functions.
* 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>
* Implement shader subgroup rotate intrinsics (#6878)Darren Wihandi2025-04-22
| | | | | | | | | | | | | * Initial implementation for SPIRV, GLSL and Metal * test add bool test * Fix and improve subgroup rotate tests * Add proper GLSL extensions and proper Metal type checking * Clean up tests and add diagnostics test for subgroup type for Metal * Update wave-intrinsics docs
* Fix quad control required SPIRV version for emit-spirv-via-glsl (#6869)Darren Wihandi2025-04-21
| | | Co-authored-by: Yong He <yonghe@outlook.com>
* Implement 64bit countbits intrinsic (#6433) (#6845)sricker-nvidia2025-04-19
| | | | | | | | | | | | | Change modifies the countbits intrinsic to use generics in order to support 64bit countbits on select platforms where this is supported. On platforms where this is not natively supported, we emulate by converting the 64-bit type into a uint2 (metal and spir-v). This should align with the implementation of other uint64_t intrinsics such as abs, min, max and clamp. Added new countbits64 test to verify changes. Updated documentation for 64bit-type-support.html
* 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.
* Implement subgroup quad operations for Metal (#6745)Darren Wihandi2025-04-04
|
* Add sparse texture Load intrinsic for SPIRV (#6702)DarrelFW3212025-04-03
| | | | | | | | | * Implement sparse texture Load intrinsics for SPIRV * changed test name from TEST_load to TEST_sparse --------- Co-authored-by: Darren Wihandi <65404740+fairywreath@users.noreply.github.com>
* Use coopvec supporting dxcompiler.dll and dxil.dll (#6719)Jay Kwak2025-04-01
| | | | | * Use coopvec supporting dxcompiler.dll and dxil.dll * Fix the failing tests
* Fix SPV_KHR_maximal_reconvergence extension name spelling (#6687)Pavel Asyutchenko2025-03-26
| | | | | | | * Fix SPV_KHR_maximal_reconvergence extension name spelling Vulkan validation layers emit warnings on lowercase khr. * Move OpExtension check
* Implement floating-point pack/unpack intrinsics for all targets (#6503)Darren Wihandi2025-03-18
| | | | | | | * Implement floating-point pack/unpack intrinsics * remove unused functions and update caps in glsl meta file * rename pack capability
* 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.
* Implement sparse texture Sample* intrinsics for SPIRV (#6377)Darren Wihandi2025-02-28
| | | | | | | | | | | | | | | | | | | | | * implement sparse residency samples for spirv * udpate test * separate tests to non-combined and combined sampler * remove expected failure * add expected failure for dx12 combined sampler test * remove expected failure * fix submodule merge * add back dx12 test failure --------- Co-authored-by: Yong He <yonghe@outlook.com>
* Add Slang-specific intrinsics for integer pack/unpack (#6459)Darren Wihandi2025-02-28
| | | | | | | | | | | | | | | | | | | | | * update hlsl meta * update test * use slang syntax in meta file * improve meta file * fix pack clamp u8 * remove builtin packed types, use typealias instead * fix wgsl pack clamp * fix formatting --------- Co-authored-by: Yong He <yonghe@outlook.com>
* Add WaveGetLane* support for Metal and WGSL (#6371)Darren Wihandi2025-02-28
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | * support WaveGetLane* for WGSL and Metal * update test and glsl support * address review comments and fix metal test * add missing pragma guard * update test * Revert "update test" This reverts commit f2b97e91c29de154190710580c343bd0764aedbb. * update failing glsl metal test and added new test * make hlsl and glsl outputs similar * update test * disable tests for Metal and cleanup * comment fix * add expected failures * correct expected failures list * remove expected failure * add tests to expected failure --------- Co-authored-by: Yong He <yonghe@outlook.com>
* Add inner texture type to reflection json (#6416)Devon2025-02-27
| | | | | | | | | | | | | * Add inner texture type to reflection json * Add expected result of test * Adjust test expected results * Fix ci test result --------- Co-authored-by: Yong He <yonghe@outlook.com>
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 17:49:31 +0000