| Age | Commit message (Collapse) | Author |
|---|
|
|
|
I think the commit diff speaks for itself.
|
|
It turned out that SPIRV-Headers update has to be done after building
SPIRV-Tools
Co-authored-by: Gangzheng Tong <tonggangzheng@gmail.com>
|
|
Before:
- Uses `LOAD_LIBRARY_SEARCH_USER_DIRS` in `LoadLibraryExW`, which might
cause exception if there is no pathes added by `AddDllDirectory()`
After:
- Use the composite flag `LOAD_LIBRARY_SEARCH_DEFAULT_DIRS`, which
searches for several locations.
- Will still search dir added by `AddDllDirectory()`, but avoids empty
path seraching if there is no AddDllDirectory() calls.
Related to https://github.com/shader-slang/slang/issues/8462
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
This allows doc only changes to PASS the CI by skipping jobs.
|
|
Remove SYSTEM flag from SPIRV-Headers to fix MacOS header precedence.
When the path is registered as SYSTEM, it is used with `-isystem` option
not `-I` option and it gets less searching order on MacOS.
When spirv.h is installed on the system directory, it will end up using
the system installed spirv.h, which is most likely an older version than
we should use.
|
|
fixes https://github.com/shader-slang/slang/issues/8472
Fixes an issue with GCC 9.4.0 on Ubuntu 20.04, it will throw an error
about PATH_MAX not being declared.
Co-authored-by: Mukund Keshava <mkeshava@nvidia.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
|
|
Add WASM FS module support for slang-playground
This change adds the necessary Emscripten build flags to export the
FileSystem (FS) module interface in the slang-wasm build:
- Adds -sMODULARIZE=1 to enable modular builds
- Adds -sEXPORTED_RUNTIME_METHODS=['FS'] to export the FS interface
These changes are required to support the slang-playground.
The existing flags are also reformatted for better readability.
Related to https://github.com/shader-slang/slang-playground/issues/170
|
|
This PR resolves a configuration conflict where the
SLANG_ENABLE_SPIRV_TOOLS_MIMALLOC option was defined twice with
different defaults, causing user settings to be overridden.
Changes:
- Removed duplicate definition in external/CMakeLists.txt that only
applied to WIN32
- Define default in main CMakeLists.txt
- ON by default for Windows
- OFF by default for other platforms (still need to resolve crashes so
set it to `OFF` for now)
Preserved user configuration - explicit
-DSLANG_ENABLE_SPIRV_TOOLS_MIMALLOC=value settings are now properly
respected
Related to https://github.com/shader-slang/slang/issues/8158
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
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
|
|
- Remove hard-coded Win SDK version.
- Using self-hosted machine for building release package in Windows.
- Remove the steps from release.yml that have been done in common steup.
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Enable GetLssPositionsAndRadii() call in rayGenLssIntrinsicsHitObject
shader that was previously commented out. This fixes the failing
ray-tracing-lss-intrinsics-hit-object test which was returning all zero
values for LSS position and radius data.
The hit object LSS intrinsics are now working correctly in D3D12
backend, returning proper endcap positions and radii values as expected
by the test. All 27 test assertions now pass successfully.
Fixes #8128
|
|
Fixes #8406 (and #8410).
`AddressSpace`, `MemoryScope` and `AccessQualifier` are no longer
`BaseType`.
I added a new `__magic_enum` (very similar to `__magic_type`) syntax to
be able to easily create values or these enums from the compiler. (I
don't know if it was the right way to do it, but it works and the
changes are small enough?).
I had a weird bug: `tests/language-feature/capability/address-of.slang`
was failing in `IRBuilder::_findOrEmitConstant(IRConstant& keyInst)`.
When needing a new `u64(0)`, it did not find it in the `ConstantMap`
first, but then failed to add it right after because it already existed
in the map! But this was triggered by `IRPtrType*
IRBuilder::getPtrType(IROp op, IRType* valueType, AccessQualifier
accessQualifier, AddressSpace addressSpace)`, which is a strange
coincidence... but I could not find the issue in what I did. I ended up
bumping unordered_dense, and it solved the issue (so there was a bug in
there).
|
|
|
|
requirements (#7269)
## Summary
This PR enhances constexpr validation by adding proper error checking
when function arguments cannot satisfy constexpr parameter requirements,
addressing issue #6370.
## Problem
Previously, when a function declared constexpr parameters, the compiler
would attempt to propagate constexpr-ness to the call site arguments,
but there was insufficient validation and error reporting when this
propagation failed. This could lead silent failures where constexpr
requirements weren't properly enforced
## Solution
This PR adds checks that:
1. **Validates constexpr arguments**: When a function parameter is
marked as `constexpr`, the compiler now explicitly checks that the
corresponding argument can be marked as `constexpr`
2. **Issues clear compilation errors**: added
`Diagnostics::argIsNotConstexpr`)
3. **Handles both call scenarios**: The validation works for both:
- Direct function calls with IR-level function definitions
- Calls to function from external modules
Fixes #6370
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
* Added optional mimalloc integration for SPIRV-Tools to improve
compilation performance
* Cloning `mimalloc` repo to external/ on demand during cmake config
* Enabled by default on Windows, configurable via
`SLANG_ENABLE_SPIRV_TOOLS_MIMALLOC` CMake option, to enable static
mimalloc build in spirv-tools.
* There are some crashes in Linux and Mac with static mimalloc enabled,
likely due to the system malloc and mimalloc mixed usage. This might be
expected as mimalloc in spirv-tools is not extensively tested according
to
https://github.com/KhronosGroup/SPIRV-Tools?tab=readme-ov-file#dependency-on-mimalloc.
So by default only Windows has this enabled.
Close: https://github.com/shader-slang/slang/issues/8158
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
The `SpirvInstructionHelper::loadBlob()` method could segfault when
calling `m_headerWords.addRange()` if the SPIR-V blob contained
insufficient data for the required 5-word header.
To reproduce, run
```
./build/Debug/bin/slangc.exe tests/modules/environment.slang -o tests/modules/environment.slang-module -target spirv -separate-debug-info
(0): error 57004: output SPIR-V contains no exported symbols. Please make sure to specify at least one entrypoint.
Segmentation fault
```
The error is expected, but the `Segmentation fault` is not.
This PR adds the check to ensure the SPIR-V blob has at least
`SPV_INDEX_INSTRUCTION_START * sizeof(SpvWord)` bytes (20 bytes minimum)
before attempting to process the header words.
Related to: https://github.com/shader-slang/slang/issues/7547
|
|
Add `findModifier` for `DeclReflection` so pattern like `extern struct
foo;` can be properly reflected.
Closes #8009
|
|
const decl. (#8392)
Closes #8184.
We fixed three issues with this regression test:
1. After generating IR for a `SpecializeComponentType`, we should also
strip the frontend
decorations from the IR so there is no HighLevelDeclDecoration that will
go into the backend.
2. When lowering a static const inside a generic function, we should not
give the static const
a linkage, because it won't such constant will not appear in global
scope. Trying to give it a
linkage decoration will lead to the parent generic (for the function) to
have two duplicate
Export/Import decorations with different mangle names, and confuses the
linker.
3. Make sure internal exceptions does not leak through
`IComponentType::getEntryPointCode`/`getTargetCode`.
|
|
Closes #8409, but ended up being more about fixing another bug. While
the issue itself seems to only be a simple typo fix (see second commit
in this PR), I found out during writing a test that pointers never got
correct locations regardless of layout. Their locations were always
assigned to zero due to lacking a resource usage entry in `TypeLayout`.
They were also missing the `Flat` decoration, so I went ahead and added
that too.
I can split this up into two separate PRs if that's preferred; both
aspects just share a test right now and fix a similar-looking issue in
the resulting SPIR-V.
|
|
|
|
…function-parameter.slang #8315
Root Cause:
CUDA compilation crashed with `assert failure:
!seenFinalUnsizedArrayField` because unsized arrays like
`RWStructuredBuffer<uint> globalBuffer[]` were not the final field in
generated parameter structs, violating the layout constraint in
slang-ir-layout.cpp.
Fix:
Extended `collectGlobalUniformParameters` to automatically reorder
struct fields for CUDA targets - regular fields first, unsized arrays
last. Other targets preserve original order.
Impact:
- Enables CUDA support for nonuniform resource indexing as function
parameters
- Zero impact on existing GLSL/HLSL/SPIRV targets
- Automatic handling - no manual parameter reordering required
Files: slang-emit.cpp, slang-ir-collect-global-uniforms.cpp/.h, test
file
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
|
## Problem
When generic functions with debug variables were specialized with
concrete types containing non-debuggable fields (e.g.,
`StructuredBuffer`), the IR cloning process would create invalid
`DebugVar` instructions without checking if the substituted types
remained debuggable.
## Solution
This fix adds a defensive check in the legalization pass that removes
the debugVar created for the non-debuggable types.
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Fixes this regression:
```slang
struct MyType
{
// Regression Condition 1: there must be more than one member in the lookup scope.
float v;
int getSum() { return 0; }
}
void m(MyType t)
{
// Regression condition 2: the completion must be in an init expression.
// Regression condition 3: none of the candidate members can coerce to the expected type.
// Regression behavior: no completion candidates are shown, because
// SemanticsVisitor::resolveOverloadedLookup throws an error when there are 0 applicable candidates
// after type coercion filtering.
Texture2D x = t.; // completion request after . here
}
```
The root cause is that we shouldn't be applying candidate filtering on
the candidate list when in completion checking mode.
Closes #8417.
|
|
This updates SPIRV submodules.
Unfortunately we cannot use the latest because one of PRs in SPIRV-Tools
causes VVL errors on our slang-tests. The issue is tracked in
- https://github.com/KhronosGroup/SPIRV-Tools/issues/6316
- https://github.com/shader-slang/slang/issues/8414
This PR also updates/simplifies the document of updating SPIRV process.
|
|
|
|
- Enable ccache in slang bulid for CTS test
- Allow manual CI dispatch
|
|
This fixes an issue where non-raytracing kernels couldn't contain any
RaytracingAccelerationStructure resources even when not used.
|
|
macos 15.6 includes python 3.9.6 with Xcode, which doesn't understand
match/case. Changing it to to the less spiffy if/elif.
Co-authored-by: Yong He <yonghe@outlook.com>
Co-authored-by: Sam Estep <sam@samestep.com>
|
|
The test can be enabled
|
|
This PR modernizes the Windows-specific code by replacing ANSI Windows
API functions with their Unicode (wide character) counterparts. This
change ensures proper handling of Unicode file paths and strings on
Windows systems.
### File Operations (`source/core/slang-io.cpp`)
- `DeleteFileA` → `DeleteFileW`
- `GetTempPathA` → `GetTempPathW`
- `GetTempFileNameA` → `GetTempFileNameW`
- `RemoveDirectoryA` → `RemoveDirectoryW`
- `SHFileOperationA` → `SHFileOperationW`
- `GetModuleFileNameA` → `GetModuleFileNameW` with UTF-8 conversion
### Platform Operations (`source/core/slang-platform.cpp`)
- `GetModuleHandleExA` → `GetModuleHandleExW`
- `LoadLibraryExA` → `LoadLibraryExW`
- `LoadLibraryA` → `LoadLibraryW`
- `OutputDebugStringA` → `OutputDebugStringW`
### Runtime and Tools
- `MessageBoxA` → `MessageBoxW` in slang-rt
- `GetCurrentDirectoryA` → `GetCurrentDirectoryW` in slang-fiddle
- String literal conversion to wide strings in vk-pipeline-create
---------
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: Gangzheng Tong <gtong-nv@users.noreply.github.com>
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Fixes #8396 by not emitting the `ArrayStride` when it would've been
zero. The problem is caused by #7848, more details in the issue
description.
I checked that with equivalent GLSL code, glslangValidator does not emit
any `ArrayStride`. I assume that the addition of `storageClass ==
SpvStorageClassStorageBuffer` to line 1848 is not a mistake. If it is,
that addition could also be simply reverted to fix this issue, I tested
that option as well. With these changes, Slang tests work locally on my
PC again.
Related to this; it'd be nice to have GPUs from multiple vendors in the
CI to avoid this kind of thing happening again. Or even just llvmpipe;
that doesn't even require a GPU and would add at least one more driver
to test with.
|
|
1. Adds retry logic - Implements 3-attempt retry mechanism for
intermittent test failures
2. Reduces parallelism for Linux - Changes server counts back to 1 for
more stable execution, given the test is short in Linux for now
3. Adds detailed error logging - Enhanced diagnostic information for
test parsing failures
4. Add Python 3.10 setup for slangpy tests for github runners
5. Removes submodule checkout for slang and slasngpy tests
6. Adds check-ci job - Implements consolidated CI status for simplified
branch protection rule
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
This change relaxes a previous restriction on link-time types and
constants, so that we now allow them to be used to define shader
parameters.
Doing so will result in a parameter layout that is incomplete prior to
linking. The PR added a test to call the reflection API on a fully
linked program and ensure that we can report correct binding info.
|
|
- This PR removes "wasm" check from the test job, since there is no test
job for "wasm" yet.
- Also, move the check for slang-rhi and slangpy test to the job level
to skip the setup etc. if no run is needed.
- Update the macos compiler version in building.md to match the clang
compiler used in CI.
- Reduce the server count in Linux and Mac when running slang-test to
ease the machine load
- Run slangpy test with `-n auto --maxprocesses=4` to speed up
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
The header execinfo.h and the related backtrace functionality is not
available on all linux platforms. In particular it's missing on musl
linux and on Android before API version 33. This causes compilation
errors on those platforms.
With this change, we first check if backtrace functionality is available
by checking if we are using glibc or a compatible Android version.
Tested on manylinux_2_28 with glibc 2.28 and musllinux_1_2 with musl
1.2, has not been tested on Android.
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
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.
|
|
The code int x4 = 0xFFFFFFFFFFFFFFFF previously did not produce a
warning due to the value being too large for the type. This patch now
checks for this and similar issues during parsing.
|
|
If a given library cannot be found using LoadLibraryExA then try again
using LoadLibraryA. Return an error only if both of these failed.
|
|
Previously we are missing `ccache --set-config=cache_dir="$ccache_dir"`
and the build was using the default cache_dir.
This PR fixed that and the cache are hit in reruns.
```
📊 ccache statistics (post-build):
Cacheable calls: 1198 / 1198 (100.0%)
Hits: 1195 / 1198 (99.75%)
Direct: 1193 / 1195 (99.83%)
Preprocessed: 2 / 1195 ( 0.17%)
Misses: 3 / 1198 ( 0.25%)
Local storage:
Cache size (GiB): 0.6 / 5.0 (11.11%)
Hits: 1195 / 1198 (99.75%)
Misses: 3 / 1198 ( 0.25%)
```
|
|
The CI has been re-organized with the following:
```
ci.yml (Main Orchestrator)
├── filter job
│ ├── Documentation Only? → Yes → Skip CI
│ └── Documentation Only? → No → Continue CI
│
├── Build Jobs
│ └── ci-slang-build.yml
│ ├── common-setup action
│ ├── Build & Package
│ └── Upload Artifacts (Build package and Tests package)
│
└── Test Jobs
└── ci-slang-test.yml
└── common-test-setup action
├── Download Tests Artifacts ← (from Build)
└── Run Tests
```
To achieve fine-grained build->test dependency, instead of using `matrix
strategy` in single build (or single test) job, the main ci.yml
statically defines the each config of the build and test job.
e.g. `build-windows-debug-cl-x86_64-gpu` and
`test-windows-debug-cl-x86_64-gpu` are a pair of the build-test job for
the windows/debug/ci config.
Closes: https://github.com/shader-slang/slang/issues/6728
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Update intrinsics signature as per the nvapi header
|
|
Enable CUDA support for additional HLSL intrinsic tests by implementing
missing functionality and fixing compiler bugs affecting CUDA targets.
- Fix critical bug in InterlockedCompareStore64 where division used /4
instead of /8 for 64-bit types, causing incorrect memory addressing for
all signed int 64_t atomics
- Add signed int64_t atomic wrappers (atomicExch, atomicCAS) to CUDA
prelu de that properly cast to/from unsigned types as required by CUDA's
atomic API
- Enable tests: atomic-intrinsics-64bit.slang
- Implement CUDA support for QuadAny and QuadAll operations using warp
shu ffle primitives (__shfl_sync with quad-level lane masking)
- Add CUDA to quad_control capability definition in
slang-capabilities.capdef
- Add _slang_quadAny/_slang_quadAll helper functions to CUDA prelude
- Enable tests: quad-control-comp-functionality.slang,
subgroup-quad.slang
---------
Co-authored-by: szihs <675653+szihs@users.noreply.github.com>
|
|
Before this PR only the following was a valid line without any
white-space character nor additional `/` character,
```
//TEST:
```
This PR is to allow slang-test to handle the following variants of the
test command comments,
```
///TEST:
// TEST:
// TEST:
////// TEST:
```
This PR revealed a regression on two tests:
- tests/cpp-compiler/c-compile-shared-library.c (cpu)
- tests/cpp-compiler/cpp-compile-shared-library.cpp (cpu)
They are disabled as a part of this PR.
And there is a new github issue to track it later,
- https://github.com/shader-slang/slang/issues/8362
|
|
closes https://github.com/shader-slang/slang/issues/3313
|
|
|
|
Add 29 intrinsics to the list by new test
|
|
|
